/* Copyright (c) 2008 - 2024 Advanced Micro Devices, Inc.

 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
 in the Software without restriction, including without limitation the rights
 to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
 copies of the Software, and to permit persons to whom the Software is
 furnished to do so, subject to the following conditions:

 The above copyright notice and this permission notice shall be included in
 all copies or substantial portions of the Software.

 THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
 IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
 FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
 AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
 LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
 OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE. */

#ifndef WITHOUT_HSA_BACKEND

#include "platform/program.hpp"
#include "platform/kernel.hpp"
#include "os/os.hpp"
#include "utils/debug.hpp"
#include "utils/flags.hpp"
#include "utils/options.hpp"
#include "utils/versions.hpp"
#include "thread/monitor.hpp"
#include "CL/cl_ext.h"

#include "vdi_common.hpp"
#include "device/comgrctx.hpp"
#include "device/devhostcall.hpp"
#include "device/rocm/rocdevice.hpp"
#include "device/rocm/rocblit.hpp"
#include "device/rocm/rocvirtual.hpp"
#include "device/rocm/rocprogram.hpp"
#include "device/rocm/rocmemory.hpp"
#include "device/rocm/rocglinterop.hpp"
#include "device/rocm/rocsignal.hpp"
#include "platform/sampler.hpp"

#if defined(__clang__)
#if __has_feature(address_sanitizer)
#include "device/rocm/rocurilocator.hpp"
#endif
#endif

#include <algorithm>
#include <cstring>
#include <fstream>
#include <iostream>
#include <memory>
#ifdef ROCCLR_SUPPORT_NUMA_POLICY
#include <numa.h>
#include <numaif.h>
#endif // ROCCLR_SUPPORT_NUMA_POLICY
#include <sstream>
#include <vector>
#endif // WITHOUT_HSA_BaCKEND

#define OPENCL_VERSION_STR XSTR(OPENCL_MAJOR) "." XSTR(OPENCL_MINOR)
#define OPENCL_C_VERSION_STR XSTR(OPENCL_C_MAJOR) "." XSTR(OPENCL_C_MINOR)


static_assert(static_cast<uint32_t>(amd::Device::VmmAccess::kNone)
              == static_cast<uint32_t>(HSA_ACCESS_PERMISSION_NONE),
              "Vmm Access Flag None mismatch with ROC-runtime!");
static_assert(static_cast<uint32_t>(amd::Device::VmmAccess::kReadOnly)
              == static_cast<uint32_t>(HSA_ACCESS_PERMISSION_RO),
              "Vmm Access Flag Read mismatch with ROCr-runtime!");
static_assert(static_cast<uint32_t>(amd::Device::VmmAccess::kWriteOnly)
              == static_cast<uint32_t>(HSA_ACCESS_PERMISSION_WO),
              "Vmm Access Flag Write mismatch with ROC-runtime!");
static_assert(static_cast<uint32_t>(amd::Device::VmmAccess::kReadWrite)
              == static_cast<uint32_t>(HSA_ACCESS_PERMISSION_RW),
              "Vmm Access Flag Read Write mismatch with ROC-runtime!");

#ifndef WITHOUT_HSA_BACKEND

namespace {

inline bool getIsaMeta(std::string isaName, amd_comgr_metadata_node_t& isaMeta) {
  amd_comgr_status_t status;
  status = amd::Comgr::get_isa_metadata(isaName.c_str(), &isaMeta);
  return (status == AMD_COMGR_STATUS_SUCCESS) ? true : false;
}

inline bool releaseIsaMeta(amd_comgr_metadata_node_t& isaMeta) {
  return AMD_COMGR_STATUS_SUCCESS == amd::Comgr::destroy_metadata(isaMeta);
}

bool getValueFromIsaMeta(amd_comgr_metadata_node_t& isaMeta, const char* key,
                         std::string& retValue) {
  amd_comgr_status_t status;
  amd_comgr_metadata_node_t valMeta;
  size_t size = 0;

  status = amd::Comgr::metadata_lookup(isaMeta, key, &valMeta);
  if (status == AMD_COMGR_STATUS_SUCCESS) {
    status = amd::Comgr::get_metadata_string(valMeta, &size, NULL);
  }
  if (status == AMD_COMGR_STATUS_SUCCESS) {
    retValue.resize(size - 1);
    status = amd::Comgr::get_metadata_string(valMeta, &size, &(retValue[0]));
  }
  if (status == AMD_COMGR_STATUS_SUCCESS) {
    status = amd::Comgr::destroy_metadata(valMeta);
  }

  return (status == AMD_COMGR_STATUS_SUCCESS) ? true : false;
}

} // namespace

namespace amd::device {
extern const char* HipExtraSourceCode;
extern const char* HipExtraSourceCodeNoGWS;
} // namespace amd::device

namespace amd::roc {
bool roc::Device::isHsaInitialized_ = false;
std::vector<hsa_agent_t> roc::Device::gpu_agents_;
std::vector<AgentInfo> roc::Device::cpu_agents_;

address Device::mg_sync_ = nullptr;

bool NullDevice::create(const amd::Isa &isa) {
  if (!isa.runtimeRocSupported()) {
    LogPrintfError("Offline HSA device %s is not supported", isa.targetId());
    return false;
  }

  online_ = false;
  // Mark the device as GPU type
  info_.type_ = CL_DEVICE_TYPE_GPU;
  info_.vendorId_ = 0x1002;

  roc::Settings* hsaSettings = new roc::Settings();
  settings_ = hsaSettings;
  if (!hsaSettings ||
      !hsaSettings->create(false, isa,
                           isa.xnack() == amd::Isa::Feature::Enabled)) {
    LogPrintfError("Error creating settings for offline HSA device %s", isa.targetId());
    return false;
  }

  if (!ValidateComgr()) {
    LogPrintfError("Code object manager initialization failed for offline HSA device %s",
                   isa.targetId());
    return false;
  }

  if (!amd::Device::create(isa)) {
    LogPrintfError("Unable to setup offline HSA device %s", isa.targetId());
    return false;
  }

  // Report the device name
  ::strncpy(info_.name_, isa.targetId(), sizeof(info_.name_) - 1);
  info_.extensions_ = getExtensionString();
  info_.maxWorkGroupSize_ = hsaSettings->maxWorkGroupSize_;
  ::strncpy(info_.vendor_, "Advanced Micro Devices, Inc.", sizeof(info_.vendor_) - 1);
  info_.oclcVersion_ = "OpenCL C " OPENCL_C_VERSION_STR " ";
  info_.spirVersions_ = "";
  std::stringstream ss;
  ss << AMD_BUILD_STRING " (HSA," << (settings().useLightning_ ? "LC" : "HSAIL");
  ss <<  ") [Offline]";
  ::strncpy(info_.driverVersion_, ss.str().c_str(), sizeof(info_.driverVersion_) - 1);
  info_.version_ = "OpenCL " OPENCL_VERSION_STR " ";
  return true;
}

Device::Device(hsa_agent_t bkendDevice)
    : mapCacheOps_(nullptr)
    , mapCache_(nullptr)
    , bkendDevice_(bkendDevice)
    , pciDeviceId_(0)
    , gpuvm_segment_max_alloc_(0)
    , alloc_granularity_(0)
    , xferQueue_(nullptr)
    , xferRead_(nullptr)
    , xferWrite_(nullptr)
    , freeMem_(0)
    , vgpusAccess_("Virtual GPU List Ops Lock", true)
    , hsa_exclusive_gpu_access_(false)
    , queuePool_(QueuePriority::Total)
    , coopHostcallBuffer_(nullptr)
    , queueWithCUMaskPool_(QueuePriority::Total)
    , numOfVgpus_(0)
    , preferred_numa_node_(0)
    , maxSdmaReadMask_(0)
    , maxSdmaWriteMask_(0) {
  group_segment_.handle = 0;
  system_segment_.handle = 0;
  system_coarse_segment_.handle = 0;
  system_kernarg_segment_.handle = 0;
  gpuvm_segment_.handle = 0;
  gpu_fine_grained_segment_.handle = 0;
  gpu_ext_fine_grained_segment_.handle = 0;
  prefetch_signal_.handle = 0;
  isXgmi_ = false;
  cache_state_ = Device::CacheState::kCacheStateInvalid;
}

void Device::setupCpuAgent() {
  int32_t numaDistance = std::numeric_limits<int32_t>::max();
  uint32_t index = 0; // 0 as default
  auto size = cpu_agents_.size();
  for (uint32_t i = 0; i < size; i++) {
    std::vector<amd::Device::LinkAttrType> link_attrs;
    link_attrs.push_back(std::make_pair(LinkAttribute::kLinkDistance, 0));
    if (findLinkInfo(cpu_agents_[i].fine_grain_pool, &link_attrs)) {
      if (link_attrs[0].second < numaDistance) {
        numaDistance = link_attrs[0].second;
        index = i;
      }
    }
  }
  std::vector<amd::Device::LinkAttrType> link_attrs;
  link_attrs.push_back(std::make_pair(LinkAttribute::kLinkLinkType, 0));
  if (findLinkInfo(cpu_agents_[0].fine_grain_pool, &link_attrs)) {
    isXgmi_ = (link_attrs[0].second == HSA_AMD_LINK_INFO_TYPE_XGMI);
  }

  preferred_numa_node_ = index;
  cpu_agent_ = cpu_agents_[index].agent;
  system_segment_ = cpu_agents_[index].fine_grain_pool;
  system_coarse_segment_ = cpu_agents_[index].coarse_grain_pool;
  system_kernarg_segment_ = cpu_agents_[index].kern_arg_pool;
  ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Numa selects cpu agent[%zu]=0x%zx(fine=0x%zx,"
          "coarse=0x%zx) for gpu agent=0x%zx CPU<->GPU XGMI=%d", index, cpu_agent_.handle,
          system_segment_.handle, system_coarse_segment_.handle, bkendDevice_.handle, isXgmi_);
}

void Device::checkAtomicSupport() {
  std::vector<amd::Device::LinkAttrType> link_attrs;
  link_attrs.push_back(std::make_pair(LinkAttribute::kLinkAtomicSupport, 0));
  if (findLinkInfo(system_segment_, &link_attrs)) {
    if (link_attrs[0].second == 1) {
      info_.pcie_atomics_ = true;
    }
  }
}

Device::~Device() {
  if (coopHostcallBuffer_) {
    amd::disableHostcalls(coopHostcallBuffer_);
    context().svmFree(coopHostcallBuffer_);
    coopHostcallBuffer_ = nullptr;
  }
  // Release cached map targets
  for (uint i = 0; mapCache_ != nullptr && i < mapCache_->size(); ++i) {
    if ((*mapCache_)[i] != nullptr) {
      (*mapCache_)[i]->release();
    }
  }
  delete mapCache_;
  delete mapCacheOps_;

  if (nullptr != p2p_stage_) {
    p2p_stage_->release();
    p2p_stage_ = nullptr;
  }
  if (nullptr != mg_sync_) {
    GlbCtx().svmFree(mg_sync_);
    mg_sync_ = nullptr;
  }
  if (glb_ctx_ != nullptr) {
      glb_ctx_->release();
      glb_ctx_ = nullptr;
  }

  for (auto& it : queuePool_) {
    for (auto qIter = it.begin(); qIter != it.end(); ) {
      hsa_queue_t* queue = qIter->first;
      auto& qInfo = qIter->second;
      if (qInfo.hostcallBuffer_) {
        ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hostcall buffer %p for hardware queue %p",
                qInfo.hostcallBuffer_, qIter->first->base_address);
        amd::disableHostcalls(qInfo.hostcallBuffer_);
        context().svmFree(qInfo.hostcallBuffer_);
      }
      ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hardware queue %p with refCount 0",
              queue->base_address);
      qIter = it.erase(qIter);
      hsa_queue_destroy(queue);
    }
  }
  queuePool_.clear();

  // Destroy temporary buffers for read/write
  delete xferRead_;
  delete xferWrite_;

  // Destroy transfer queue
  delete xferQueue_;

  delete blitProgram_;

  if (context_ != nullptr) {
    context_->release();
  }

  delete[] p2p_agents_list_;

  if (0 != prefetch_signal_.handle) {
    hsa_signal_destroy(prefetch_signal_);
  }
}

bool NullDevice::initCompiler(bool isOffline) {
  return true;
}

bool NullDevice::destroyCompiler() {
  return true;
}

void NullDevice::tearDown() { destroyCompiler(); }

bool NullDevice::init() {
  // Initialize the compiler
  if (!initCompiler(offlineDevice_)) {
    return false;
  }

  // Create offline devices for all ISAs not already associated with an online
  // device. This allows code objects to be compiled for all supported ISAs.
  std::vector<Device*> devices = getDevices(CL_DEVICE_TYPE_GPU, false);
  for (const amd::Isa *isa = amd::Isa::begin(); isa != amd::Isa::end(); isa++) {
    if (!isa->runtimeRocSupported()) {
      continue;
    }
    bool isOnline = false;
    // Check if the particular device is online
    for (size_t i = 0; i < devices.size(); i++) {
      if (&(devices[i]->isa()) == isa) {
        isOnline = true;
        break;
      }
    }
    if (isOnline) {
      continue;
    }
    std::unique_ptr<NullDevice> nullDevice(new NullDevice());
    if (!nullDevice) {
      LogPrintfError("Error allocating new instance of offline HSA device %s", isa->targetId());
      return false;
    }
    if (!nullDevice->create(*isa)) {
      LogPrintfError("Skipping creating new instance of offline HSA sevice %s", isa->targetId());
      continue;
    }
    nullDevice.release()->registerDevice();
  }
  return true;
}

NullDevice::~NullDevice() {
}

hsa_status_t Device::iterateAgentCallback(hsa_agent_t agent, void* data) {
  hsa_device_type_t dev_type = HSA_DEVICE_TYPE_CPU;

  hsa_status_t stat = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev_type);

  if (stat != HSA_STATUS_SUCCESS) {
    LogPrintfError("HSA_AGENT_INFO_DEVICE failed with %x", stat);
    return stat;
  }

  if (dev_type == HSA_DEVICE_TYPE_CPU) {
    AgentInfo info = { agent, { 0 }, { 0 }, { 0 }};
    stat = hsa_amd_agent_iterate_memory_pools(agent, Device::iterateCpuMemoryPoolCallback,
                                              reinterpret_cast<void*>(&info));
    if (stat == HSA_STATUS_SUCCESS) {
      cpu_agents_.push_back(info);
    }
  } else if (dev_type == HSA_DEVICE_TYPE_GPU) {
    gpu_agents_.push_back(agent);
  }

  return stat;
}

hsa_ven_amd_loader_1_00_pfn_t Device::amd_loader_ext_table = {nullptr};

hsa_status_t Device::loaderQueryHostAddress(const void* device, const void** host) {
  return amd_loader_ext_table.hsa_ven_amd_loader_query_host_address
      ? amd_loader_ext_table.hsa_ven_amd_loader_query_host_address(device, host)
      : HSA_STATUS_ERROR;
}

Device::XferBuffers::~XferBuffers() {
  // Destroy temporary buffer for reads
  for (const auto& buf : freeBuffers_) {
    delete buf;
  }
  freeBuffers_.clear();
}

bool Device::XferBuffers::create() {
  Memory* xferBuf = nullptr;
  bool result = false;

  // Create a buffer object
  xferBuf = new Buffer(dev(), bufSize_);

  // Try to allocate memory for the transfer buffer
  if ((nullptr == xferBuf) || !xferBuf->create()) {
    delete xferBuf;
    xferBuf = nullptr;
    LogError("Couldn't allocate a transfer buffer!");
  } else {
    result = true;
    freeBuffers_.push_back(xferBuf);
  }

  return result;
}

Memory& Device::XferBuffers::acquire() {
  Memory* xferBuf = nullptr;
  size_t listSize;

  // Lock the operations with the staged buffer list
  amd::ScopedLock l(lock_);
  listSize = freeBuffers_.size();

  // If the list is empty, then attempt to allocate a staged buffer
  if (listSize == 0) {
    // Allocate memory
    xferBuf = new Buffer(dev(), bufSize_);

    // Allocate memory for the transfer buffer
    if ((nullptr == xferBuf) || !xferBuf->create()) {
      delete xferBuf;
      xferBuf = nullptr;
      LogError("Couldn't allocate a transfer buffer!");
    } else {
      ++acquiredCnt_;
    }
  }

  if (xferBuf == nullptr) {
    xferBuf = *(freeBuffers_.begin());
    freeBuffers_.erase(freeBuffers_.begin());
    ++acquiredCnt_;
  }

  return *xferBuf;
}

void Device::XferBuffers::release(VirtualGPU& gpu, Memory& buffer) {
  // Make sure buffer isn't busy on the current VirtualGPU, because
  // the next aquire can come from different queue
  //    buffer.wait(gpu);
  // Lock the operations with the staged buffer list
  amd::ScopedLock l(lock_);
  freeBuffers_.push_back(&buffer);
  --acquiredCnt_;
}

// ================================================================================================
bool Device::init() {
  ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Initializing HSA stack.");
  hsa_status_t status = HSA_STATUS_SUCCESS;
  // Initialize the compiler
  if (!initCompiler(offlineDevice_)) {
    LogError("initCompiler failed.");
    return false;
  }

  status = hsa_init();

  // If there are no GPUs available, hsa_init will fail with HSA_STATUS_ERROR_OUT_OF_RESOURCES
  // but for NoGpu tests to pass, true needs to be returned
  constexpr bool kNoOfflineDevices = false;
  std::vector<amd::Device*> devices = getDevices(CL_DEVICE_TYPE_GPU, kNoOfflineDevices);
  if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES && devices.size() == 0) {
    return true;
  }

  if (status != HSA_STATUS_SUCCESS) {
    LogPrintfError("hsa_init failed with %x", status);
    return false;
  }

  hsa_system_get_major_extension_table(HSA_EXTENSION_AMD_LOADER, 1, sizeof(amd_loader_ext_table),
                                       &amd_loader_ext_table);

  status = hsa_iterate_agents(iterateAgentCallback, nullptr);
  if (status != HSA_STATUS_SUCCESS) {
    LogPrintfError("hsa_iterate_agents failed with %x", status);
    return false;
  }

  std::string ordinals = amd::IS_HIP ? ((HIP_VISIBLE_DEVICES[0] != '\0') ?
                         HIP_VISIBLE_DEVICES : CUDA_VISIBLE_DEVICES)
                         : GPU_DEVICE_ORDINAL;
  if (ordinals[0] != '\0') {
    size_t pos = 0;
    std::vector<hsa_agent_t> valid_agents;
    std::set<size_t> valid_indexes;
    do {
      size_t end;
      bool deviceIdValid = true;
      end = ordinals.find_first_of(',', pos);
      if (end == std::string::npos) {
        end = ordinals.size();
      }
      std::string str_id = ordinals.substr(pos, end - pos);
      // If Uuid is specified, then convert it to index
      // Uuid is an Ascii string with a maximum of 21 chars including NULL
      // The string value is in the format GPU-<body>, <body> encodes UUID as a 16 chars hex
      if (str_id.find("GPU-") != std::string::npos) {
        for (int i = 0; i < gpu_agents_.size(); i++) {
          auto agent = gpu_agents_[i];
          char unique_id[32] = {0};
          if (HSA_STATUS_SUCCESS ==
            hsa_agent_get_info(agent, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_UUID),
                              unique_id)) {
            if (std::string(unique_id).find(str_id) != std::string::npos) {
              str_id = std::to_string(i);
              break;
            }
          }
        }
      }
      int index = atoi(str_id.c_str());
      if (index < 0 ||
          static_cast<size_t>(index) >= gpu_agents_.size() ||
          str_id != std::to_string(index)) {
        deviceIdValid = false;
      }

      if (!deviceIdValid) {
        // Exit the loop as anything to the right of invalid deviceId
        // has to be discarded
        break;
      } else {
        if (valid_indexes.find(index) == valid_indexes.end()) {
          valid_agents.push_back(gpu_agents_[index]);
          valid_indexes.insert(index);
        }
      }
      pos = end + 1;
    } while (pos < ordinals.size());
    gpu_agents_ = valid_agents;
  }

  LogPrintfInfo("Enumerated GPU agents = %lu", gpu_agents_.size());

  for (auto agent : gpu_agents_) {
    std::unique_ptr<Device> roc_device(new Device(agent));
    if (!roc_device) {
      LogError("Error creating new instance of Device on then heap.");
      continue;
    }

    if (!roc_device->create()) {
      LogError("Error creating new instance of Device.");
      continue;
    }

    // Setup System Memory to be Non-Coherent per user
    // request via environment variable. By default the
    // System Memory is setup to be Coherent
    if (roc_device->settings().enableNCMode_) {
      hsa_status_t err = hsa_amd_coherency_set_type(agent, HSA_AMD_COHERENCY_TYPE_NONCOHERENT);
      if (err != HSA_STATUS_SUCCESS) {
        LogError("Unable to set NC memory policy!");
        continue;
      }
    }

    // Check to see if a global CU mask is requested
    if (amd::IS_HIP && ROC_GLOBAL_CU_MASK[0] != '\0') {
      roc_device->getGlobalCUMask(ROC_GLOBAL_CU_MASK);
    }

    roc_device.release()->registerDevice();
  }

  // Query active devices only
  devices = getDevices(CL_DEVICE_TYPE_GPU, kNoOfflineDevices);
  if (devices.size() > 0) {
    bool p2p_available = false;
    // Loop through all available devices
    for (auto device1: devices) {
      // Find all agents that can have access to the current device
      for (auto agent: static_cast<Device*>(device1)->p2pAgents()) {
        // Find cl_device_id associated with the current agent
        for (auto device2: devices) {
          if (agent.handle == static_cast<Device*>(device2)->getBackendDevice().handle) {
            // Device2 can have access to device1
            device2->p2pDevices_.push_back(as_cl(device1));
            device1->p2p_access_devices_.push_back(device2);
            p2p_available = true;
          }
        }
      }
    }

    // Create a dummy context for internal memory allocations on all reported devices
    glb_ctx_ = new amd::Context(devices, amd::Context::Info());
    if (glb_ctx_ == nullptr) {
      LogError("glb_ctx failed");
      return false;
    }

    // Allocate a staging buffer for P2P emulation path
    if ((devices.size() >= 1) && !p2p_available) {
      amd::Buffer* buf =
          new (*glb_ctx_) amd::Buffer(*glb_ctx_, CL_MEM_ALLOC_HOST_PTR, kP2PStagingSize);
      if ((buf != nullptr) && buf->create()) {
        p2p_stage_ = buf;
      } else {
        delete buf;
        LogError("p2p stg buffer alloc failed");
        return false;
      }
    }

    // Allocate mgpu sync buffer for cooperative launches
    if (amd::IS_HIP) {
      mg_sync_ = reinterpret_cast<address>(glb_ctx_->svmAlloc(kMGInfoSizePerDevice * devices.size(),
        kMGInfoSizePerDevice, (CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS)));
      if (mg_sync_ == nullptr) {
        LogError("mgpu sync buffer alloc failed");
        return false;
      }
    }
  }

  return true;
}

extern const char* SchedulerSourceCode;

void Device::tearDown() {
  NullDevice::tearDown();
  hsa_shut_down();
}

// ================================================================================================
bool Device::create() {
  char agent_name[64] = {0};
  if (HSA_STATUS_SUCCESS != hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_NAME, agent_name)) {
    LogError("Unable to get HSA device name");
    return false;
  }

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_CHIP_ID,
                         &pciDeviceId_)) {
    LogPrintfError("Unable to get PCI ID of HSA device %s", agent_name);
    return false;
  }

  struct agent_isas_t {
    uint count;
    hsa_isa_t first_isa;
  } agent_isas = {0, {0}};
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_iterate_isas(bkendDevice_,
                             [](hsa_isa_t isa, void* data) {
                               agent_isas_t* agent_isas = static_cast<agent_isas_t*>(data);
                               if (agent_isas->count++ == 0) {
                                 agent_isas->first_isa = isa;
                               }
                               return HSA_STATUS_SUCCESS;
                             },
                             &agent_isas)) {
    LogPrintfError("Unable to iterate supported ISAs for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }
  if (agent_isas.count != 1) {
    LogPrintfError("HSA device %s (PCI ID %x) has %u ISAs but can only support a single ISA",
                   agent_name, pciDeviceId_, agent_isas.count);
    return false;
  }

  uint32_t isa_name_length = 0;
  if (HSA_STATUS_SUCCESS !=
      hsa_isa_get_info_alt(agent_isas.first_isa, (hsa_isa_info_t)HSA_ISA_INFO_NAME_LENGTH,
                           &isa_name_length)) {
    LogPrintfError("Unable to get ISA name length for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }

  std::vector<char> isa_name(isa_name_length + 1, '\0');
  if (HSA_STATUS_SUCCESS !=
      hsa_isa_get_info_alt(agent_isas.first_isa, (hsa_isa_info_t)HSA_ISA_INFO_NAME,
                           isa_name.data())) {
    LogPrintfError("Unable to get ISA name for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }

  const amd::Isa *isa = amd::Isa::findIsa(isa_name.data());
  if (!isa || !isa->runtimeRocSupported()) {
    LogPrintfError("Unsupported HSA device %s (PCI ID %x) for ISA %s", agent_name, pciDeviceId_,
                   isa_name.data());
    return false;
  }

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_PROFILE, &agent_profile_)) {
    LogPrintfError("Unable to get profile for HSA device %s (PCI ID %x)", agent_name, pciDeviceId_);
    return false;
  }

  uint32_t coop_groups = 0;
  // Check cooperative groups for HIP only
  if (amd::IS_HIP &&
      (HSA_STATUS_SUCCESS !=
       hsa_agent_get_info(bkendDevice_,
                          static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_COOPERATIVE_QUEUES),
                          &coop_groups))) {
    LogPrintfError(
        "Unable to determine if cooperative queues are supported for HSA device %s (PCI ID %x)",
        agent_name, pciDeviceId_);
    return false;
  }

  setupCpuAgent();

  // Get Agent HDP Flush Register Memory
  hsa_amd_hdp_flush_t hdpInfo;
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
        static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_HDP_FLUSH), &hdpInfo)) {
    LogPrintfError("Unable to determine HDP flush info for HSA device %s", agent_name);
    return false;
  }

  info_.hdpMemFlushCntl = hdpInfo.HDP_MEM_FLUSH_CNTL;
  info_.hdpRegFlushCntl = hdpInfo.HDP_REG_FLUSH_CNTL;
  bool hasValidHDPFlush =
      (info_.hdpMemFlushCntl != nullptr) && (info_.hdpRegFlushCntl != nullptr);

  // Create HSA settings
  assert(!settings_);
  roc::Settings* hsaSettings = new roc::Settings();
  settings_ = hsaSettings;
  if (!hsaSettings ||
      !hsaSettings->create((agent_profile_ == HSA_PROFILE_FULL), *isa,
                           isa->xnack() == amd::Isa::Feature::Enabled,
                           coop_groups, isXgmi_, hasValidHDPFlush)) {
    LogPrintfError("Unable to create settings for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }

  if (!ValidateComgr()) {
    LogPrintfError("Code object manager initialization failed for HSA device %s (PCI ID %x)",
                   agent_name, pciDeviceId_);
    return false;
  }

  if (!amd::Device::create(*isa)) {
    LogPrintfError("Unable to setup device for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }

  uint32_t hsa_bdf_id = 0;
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
        static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_BDFID), &hsa_bdf_id)) {
    LogPrintfError("Unable to determine BFD ID for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }

  info_.deviceTopology_.pcie.type = CL_DEVICE_TOPOLOGY_TYPE_PCIE_AMD;
  info_.deviceTopology_.pcie.bus = (hsa_bdf_id & (0xFF << 8)) >> 8;
  info_.deviceTopology_.pcie.device = (hsa_bdf_id & (0x1F << 3)) >> 3;
  info_.deviceTopology_.pcie.function = (hsa_bdf_id & 0x07);
  uint32_t pci_domain_id = 0;
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
        static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_DOMAIN), &pci_domain_id)) {
    LogPrintfError("Unable to determine domain ID for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }
  info_.pciDomainID = pci_domain_id;

  if (populateOCLDeviceConstants() == false) {
    LogPrintfError("populateOCLDeviceConstants failed for HSA device %s (PCI ID %x)", agent_name,
                   pciDeviceId_);
    return false;
  }
  hsaSettings->limit_blit_wg_ = info().maxComputeUnits_;
  if (!flagIsDefault(DEBUG_CLR_LIMIT_BLIT_WG)) {
    hsaSettings->limit_blit_wg_ = std::max(DEBUG_CLR_LIMIT_BLIT_WG, 0x1U);
  }
  amd::Context::Info info = {0};
  std::vector<amd::Device*> devices;
  devices.push_back(this);

  // Create a dummy context
  context_ = new amd::Context(devices, info);
  if (context_ == nullptr) {
    return false;
  }

  mapCacheOps_ = new amd::Monitor("Map Cache Lock", true);
  if (nullptr == mapCacheOps_) {
    return false;
  }

  mapCache_ = new std::vector<amd::Memory*>();
  if (mapCache_ == nullptr) {
    return false;
  }
  // Use just 1 entry by default for the map cache
  mapCache_->push_back(nullptr);

  if (settings().stagedXferSize_ != 0) {
    // Initialize staged write buffers
    if (settings().stagedXferWrite_) {
      xferWrite_ = new XferBuffers(*this, amd::alignUp(settings().stagedXferSize_, 4 * Ki));
      if ((xferWrite_ == nullptr) || !xferWrite_->create()) {
        LogError("Couldn't allocate transfer buffer objects for read");
        return false;
      }
    }

    // Initialize staged read buffers
    if (settings().stagedXferRead_) {
      xferRead_ = new XferBuffers(*this, amd::alignUp(settings().stagedXferSize_, 4 * Ki));
      if ((xferRead_ == nullptr) || !xferRead_->create()) {
        LogError("Couldn't allocate transfer buffer objects for write");
        return false;
      }
    }
  }

  // Create signal for HMM prefetch operation on device
  if (HSA_STATUS_SUCCESS != hsa_signal_create(kInitSignalValueOne, 0, nullptr, &prefetch_signal_)) {
    return false;
  }

  return true;
}

// ================================================================================================
device::Program* NullDevice::createProgram(amd::Program& owner, amd::option::Options* options) {
  device::Program* program;
  if (settings().useLightning_) {
    program = new LightningProgram(*this, owner);
  } else {
    program = new HSAILProgram(*this, owner);
  }

  if (program == nullptr) {
    LogError("Memory allocation has failed!");
  }

  return program;
}

bool Device::AcquireExclusiveGpuAccess() {
  // Lock the virtual GPU list
  vgpusAccess().lock();

  // Find all available virtual GPUs and lock them
  // from the execution of commands
  for (uint idx = 0; idx < vgpus().size(); ++idx) {
    vgpus()[idx]->execution().lock();
    // Make sure a wait is done
    vgpus()[idx]->releaseGpuMemoryFence();
  }
  if (!hsa_exclusive_gpu_access_) {
    // @todo call rocr
    hsa_exclusive_gpu_access_ = true;
  }
  return true;
}

void Device::ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const {
  // Make sure the operation is done
  vgpu.releaseGpuMemoryFence();

  // Find all available virtual GPUs and unlock them
  // for the execution of commands
  for (uint idx = 0; idx < vgpus().size(); ++idx) {
    vgpus()[idx]->execution().unlock();
  }

  // Unock the virtual GPU list
  vgpusAccess().unlock();
}

bool Device::createBlitProgram() {
  bool result = true;
  std::string extraKernel;

#if defined(USE_COMGR_LIBRARY)
  if (settings().useLightning_) {
    if (amd::IS_HIP) {
      if (settings().gwsInitSupported_) {
        extraKernel = device::HipExtraSourceCode;
      } else {
        extraKernel = device::HipExtraSourceCodeNoGWS;
      }
    } else {
      extraKernel = SchedulerSourceCode;
    }

  }
#endif  // USE_COMGR_LIBRARY

  blitProgram_ = new BlitProgram(context_);
  // Create blit programs
  if (blitProgram_ == nullptr || !blitProgram_->create(this, extraKernel, "")) {
    delete blitProgram_;
    blitProgram_ = nullptr;
    LogError("Couldn't create blit kernels!");
    return false;
  }

  return result;
}

device::Program* Device::createProgram(amd::Program& owner, amd::option::Options* options) {
  device::Program* program;
  if (settings().useLightning_) {
    program = new LightningProgram(*this, owner);
  } else {
    program = new HSAILProgram(*this, owner);
  }

  if (program == nullptr) {
    LogError("Memory allocation has failed!");
  }

  return program;
}

hsa_status_t Device::iterateGpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, void* data) {
  if (data == nullptr) {
    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
  }

  hsa_region_segment_t segment_type = (hsa_region_segment_t)0;
  hsa_status_t stat =
      hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment_type);
  if (stat != HSA_STATUS_SUCCESS) {
    return stat;
  }

  // TODO: system and device local segment
  Device* dev = reinterpret_cast<Device*>(data);
  switch (segment_type) {
    case HSA_REGION_SEGMENT_GLOBAL: {
      if (dev->settings().enableLocalMemory_) {
        uint32_t global_flag = 0;
        hsa_status_t stat =
            hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &global_flag);
        if (stat != HSA_STATUS_SUCCESS) {
          return stat;
        }

        // If the flag set is ext scoped fine grain, break the loop
        if ((global_flag & HSA_REGION_GLOBAL_FLAG_EXTENDED_SCOPE_FINE_GRAINED) != 0) {
          dev->gpu_ext_fine_grained_segment_ = pool;
          break;
        }

        if ((global_flag & HSA_REGION_GLOBAL_FLAG_FINE_GRAINED) != 0) {
          dev->gpu_fine_grained_segment_ = pool;
        } else if ((global_flag & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) != 0) {
          dev->gpuvm_segment_ = pool;

          // If cpu agent cannot access this pool, the device does not support large bar.
          hsa_amd_memory_pool_access_t tmp{};
          hsa_amd_agent_memory_pool_get_info(
            dev->cpu_agent_,
            pool,
            HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS,
            &tmp);

          if (tmp == HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED) {
            dev->info_.largeBar_ = false;
          } else {
            dev->info_.largeBar_ = ROC_ENABLE_LARGE_BAR;
          }

          // Query the recommended granularity for this pool.
          stat = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
                                              &(dev->info_.virtualMemAllocGranularity_));
          if (stat != HSA_STATUS_SUCCESS) {
            LogPrintfError("Cannot query HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE info"
                           "failed with hsa_status: %d \n", stat);
          }
        }

        if (dev->gpuvm_segment_.handle == 0) {
          dev->gpuvm_segment_ = pool;
        }
      }
      break;
    }
    case HSA_REGION_SEGMENT_GROUP:
      dev->group_segment_ = pool;
      break;
    default:
      break;
  }

  return HSA_STATUS_SUCCESS;
}

hsa_status_t Device::iterateCpuMemoryPoolCallback(hsa_amd_memory_pool_t pool, void* data) {
  if (data == nullptr) {
    LogError("CpuMemoryPoolCallback invalid args");
    return HSA_STATUS_ERROR_INVALID_ARGUMENT;
  }

  hsa_region_segment_t segment_type = (hsa_region_segment_t)0;
  hsa_status_t stat =
      hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment_type);
  if (stat != HSA_STATUS_SUCCESS) {
    LogPrintfError("HSA_AMD_MEMORY_POOL_INFO_SEGMENT query failed with %x", stat);
    return stat;
  }
  AgentInfo* agentInfo = reinterpret_cast<AgentInfo*>(data);

  switch (segment_type) {
    case HSA_REGION_SEGMENT_GLOBAL: {
      uint32_t global_flag = 0;
      stat = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS,
                                          &global_flag);
      if (stat != HSA_STATUS_SUCCESS) {
        LogPrintfError("HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS query failed with %x", stat);
        break;
      }

      // If the flag set is ext scoped fine grain, break the loop
      if ((global_flag & HSA_REGION_GLOBAL_FLAG_EXTENDED_SCOPE_FINE_GRAINED) != 0) {
        agentInfo->ext_fine_grain_pool_ = pool;
        break;
      }

      if ((global_flag & HSA_REGION_GLOBAL_FLAG_FINE_GRAINED) != 0) {
        if (agentInfo->fine_grain_pool.handle == 0) {
          agentInfo->fine_grain_pool = pool;
        } else if ((global_flag & HSA_REGION_GLOBAL_FLAG_KERNARG) == 0) {
          // If the fine_grain_pool was already filled, but kern_args flag was not set over-write.
          // This means this is region-1(fine_grain only), so over-write this with memory pool set
          // from "fine_grain and kern_args".
          agentInfo->fine_grain_pool = pool;
        }
        guarantee(((global_flag & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) == 0),
                  "Memory Segment cannot be both coarse and fine grained");
      } else {
        // If the flag is not set to fine grained, then it is coarse_grained by default.
        agentInfo->coarse_grain_pool = pool;
        guarantee(((global_flag & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) != 0),
                  "Memory Segments that are not fine grained has to be coarse grained");
        guarantee(((global_flag & HSA_REGION_GLOBAL_FLAG_FINE_GRAINED) == 0),
                  "Memory Segment cannot be both coarse and fine grained");
        guarantee(((global_flag & HSA_REGION_GLOBAL_FLAG_KERNARG) == 0),
                  "Coarse grained memory segment cannot have kern_args tag");
      }

      if ((global_flag & HSA_REGION_GLOBAL_FLAG_KERNARG) != 0) {
        agentInfo->kern_arg_pool = pool;
        guarantee(((global_flag & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) == 0),
                  "Coarse grained memory segment cannot have kern_args tag");
      }

      break;
    }
    default:
      break;
  }

  return stat;
}

bool Device::createSampler(const amd::Sampler& owner, device::Sampler** sampler) const {
  *sampler = nullptr;
  Sampler* gpuSampler = new Sampler(*this);
  if ((nullptr == gpuSampler) || !gpuSampler->create(owner)) {
    delete gpuSampler;
    return false;
  }
  *sampler = gpuSampler;
  return true;
}

void Sampler::fillSampleDescriptor(hsa_ext_sampler_descriptor_t& samplerDescriptor,
                                   const amd::Sampler& sampler) const {
  samplerDescriptor.filter_mode = sampler.filterMode() == CL_FILTER_NEAREST
      ? HSA_EXT_SAMPLER_FILTER_MODE_NEAREST
      : HSA_EXT_SAMPLER_FILTER_MODE_LINEAR;
  samplerDescriptor.coordinate_mode = sampler.normalizedCoords()
      ? HSA_EXT_SAMPLER_COORDINATE_MODE_NORMALIZED
      : HSA_EXT_SAMPLER_COORDINATE_MODE_UNNORMALIZED;
  switch (sampler.addressingMode()) {
    case CL_ADDRESS_CLAMP_TO_EDGE:
      samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_EDGE;
      break;
    case CL_ADDRESS_REPEAT:
      samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_REPEAT;
      break;
    case CL_ADDRESS_CLAMP:
      samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_CLAMP_TO_BORDER;
      break;
    case CL_ADDRESS_MIRRORED_REPEAT:
      samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_MIRRORED_REPEAT;
      break;
    case CL_ADDRESS_NONE:
      samplerDescriptor.address_mode = HSA_EXT_SAMPLER_ADDRESSING_MODE_UNDEFINED;
      break;
    default:
      return;
  }
}

bool Sampler::create(const amd::Sampler& owner) {
  hsa_ext_sampler_descriptor_t samplerDescriptor;
  fillSampleDescriptor(samplerDescriptor, owner);

  hsa_status_t status = hsa_ext_sampler_create(dev_.getBackendDevice(), &samplerDescriptor, &hsa_sampler);

  if (HSA_STATUS_SUCCESS != status) {
    DevLogPrintfError("Sampler creation failed with status: %d \n", status);
    return false;
  }

  hwSrd_ = reinterpret_cast<uint64_t>(hsa_sampler.handle);
  hwState_ = reinterpret_cast<address>(hsa_sampler.handle);

  return true;
}

Sampler::~Sampler() {
  hsa_ext_sampler_destroy(dev_.getBackendDevice(), hsa_sampler);
}

Memory* Device::getGpuMemory(amd::Memory* mem) const {
  return static_cast<roc::Memory*>(mem->getDeviceMemory(*this));
}

const bool Device::isFineGrainSupported() const {
  bool result = (info().svmCapabilities_ & CL_DEVICE_SVM_ATOMICS) != 0 ? true : false;
  if (result) {
    if (gpu_fine_grained_segment_.handle != 0) {
      return true;
    }
  }
  return false;
}
// ================================================================================================
bool Device::populateOCLDeviceConstants() {
  info_.available_ = true;

  ::strncpy(info_.name_, isa().targetId(), sizeof(info_.name_) - 1);
  char device_name[64] = {0};
  if (HSA_STATUS_SUCCESS == hsa_agent_get_info(bkendDevice_,
                                               (hsa_agent_info_t)HSA_AMD_AGENT_INFO_PRODUCT_NAME,
                                               device_name)) {
    ::strncpy(info_.boardName_, device_name, sizeof(info_.boardName_) - 1);
  }

  char unique_id[32] = {0};
  if (HSA_STATUS_SUCCESS ==
      hsa_agent_get_info(bkendDevice_, static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_UUID),
                         unique_id)) {
    // ROCr gives the UUID info in the format GPU-XXXX with length 20 bytes
    // Strip the first 4 bytes and store only the 16 bytes representing UUID
    for (size_t i = 0; i < 16; i++) {
     info_.uuid_[i] = unique_id[i+4];
    }
  }
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
                         (amd::IS_HIP) ?
                         (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COOPERATIVE_COMPUTE_UNIT_COUNT :
                         (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT,
                         &info_.maxComputeUnits_)) {
    return false;
  }
  assert(info_.maxComputeUnits_ > 0);

  info_.maxComputeUnits_ = settings().enableWgpMode_
      ? info_.maxComputeUnits_ / 2
      : info_.maxComputeUnits_;

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT,
                         &info_.maxPhysicalComputeUnits_)) {
    return false;
  }
  assert(info_.maxPhysicalComputeUnits_ > 0);

  info_.maxPhysicalComputeUnits_ = settings().enableWgpMode_
      ? info_.maxPhysicalComputeUnits_ / 2
      : info_.maxPhysicalComputeUnits_;

  if (HSA_STATUS_SUCCESS != hsa_agent_get_info(bkendDevice_,
                                               (hsa_agent_info_t)HSA_AMD_AGENT_INFO_CACHELINE_SIZE,
                                               &info_.globalMemCacheLineSize_)) {
    return false;
  }
  assert(info_.globalMemCacheLineSize_ > 0);

  // override cache line size to 256 for gfx12, as it is used for kern arg alignment.
  if ((isa().versionMajor() >= 12) && (info_.globalMemCacheLineSize_ < 256)) {
    info_.globalMemCacheLineSize_ = 256;
  }

  uint32_t cachesize[4] = {0};
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_CACHE_SIZE, cachesize)) {
    return false;
  }
  assert(cachesize[0] > 0);
  info_.globalMemCacheSize_ = cachesize[0];

  info_.globalMemCacheType_ = CL_READ_WRITE_CACHE;

  info_.type_ = CL_DEVICE_TYPE_GPU;

  info_.extensions_ = getExtensionString();
  info_.nativeVectorWidthDouble_ = info_.preferredVectorWidthDouble_ =
      (settings().doublePrecision_) ? 1 : 0;

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_CLOCK_FREQUENCY,
                         &info_.maxEngineClockFrequency_)) {
    return false;
  }

  if (!(isa().versionMajor() == 9 && isa().versionMinor() == 0 && isa().versionStepping() == 2)) {
    if (info_.maxEngineClockFrequency_ <= 0) {
      LogError("maxEngineClockFrequency_ is NOT positive!");
    }
  }

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MEMORY_MAX_FREQUENCY,
          &info_.maxMemoryClockFrequency_)) {
      return false;
  }

  uint64_t wallClockFrequency = 0; // in Hz
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY ,
          &wallClockFrequency)) {
    LogWarning("HSA_AMD_AGENT_INFO_TIMESTAMP_FREQUENCY cannot be queried. Ignored!");
  }
  info_.wallClockFrequency_ = static_cast<uint32_t>(wallClockFrequency / 1000); // in KHz

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MEMORY_WIDTH),
                         &info_.globalMemChannels_)) {
    return false;
  }
  assert(info_.globalMemChannels_ > 0);

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_DRIVER_NODE_ID),
                         &info_.driverNodeId_)) {
    return false;
  }

  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SDMA_ENG),
                         &info_.numSDMAengines_)) {
    return false;
  }

  for (uint32_t i = 0; i < info_.numSDMAengines_; i++) {
    engineAssignMap_[1 << i] = 0;
  }


  checkAtomicSupport();

  assert(system_segment_.handle != 0);
  if (HSA_STATUS_SUCCESS != hsa_amd_agent_iterate_memory_pools(
                                bkendDevice_, Device::iterateGpuMemoryPoolCallback, this)) {
    return false;
  }

  assert(group_segment_.handle != 0);

  for (auto agent: gpu_agents_) {
    if (agent.handle != bkendDevice_.handle) {
      hsa_status_t err;
      // Can another GPU (agent) have access to the current GPU memory pool (gpuvm_segment_)?
      hsa_amd_memory_pool_access_t access;
      err = hsa_amd_agent_memory_pool_get_info(agent, gpuvm_segment_,
                                               HSA_AMD_AGENT_MEMORY_POOL_INFO_ACCESS, &access);
      if (err != HSA_STATUS_SUCCESS) {
        continue;
      }

      // Find accessible p2p agents - i.e != HSA_AMD_MEMORY_POOL_ACCESS_NEVER_ALLOWED
      if (HSA_AMD_MEMORY_POOL_ACCESS_ALLOWED_BY_DEFAULT == access ||
          HSA_AMD_MEMORY_POOL_ACCESS_DISALLOWED_BY_DEFAULT == access) {
        // Agent can have access to the current gpuvm_segment_
        p2p_agents_.push_back(agent);
      }
    }
  }

  // Keep track of all P2P Agents in a Array including current device handle for IPC
  p2p_agents_list_ = new hsa_agent_t[1 + p2p_agents_.size()];
  p2p_agents_list_[0] = getBackendDevice();
  for (size_t agent_idx = 0; agent_idx < p2p_agents_.size(); ++agent_idx) {
    p2p_agents_list_[1 + agent_idx] = p2p_agents_[agent_idx];
  }

  size_t group_segment_size = 0;
  if (HSA_STATUS_SUCCESS != hsa_amd_memory_pool_get_info(group_segment_,
                                                         HSA_AMD_MEMORY_POOL_INFO_SIZE,
                                                         &group_segment_size)) {
    return false;
  }
  assert(group_segment_size > 0);

  // Find SDMA read mask
  if (HSA_STATUS_SUCCESS != hsa_amd_memory_copy_engine_status(getCpuAgent(), getBackendDevice(),
                                                              &maxSdmaReadMask_)) {
    return false;
  }
  assert(maxSdmaReadMask_ > 0 && "No SDMA engines available for Read");

  // Find SDMA write mask
  if (HSA_STATUS_SUCCESS != hsa_amd_memory_copy_engine_status(getBackendDevice(), getCpuAgent(),
                                                              &maxSdmaWriteMask_)) {
    return false;
  }
  assert(maxSdmaWriteMask_ > 0 && "No SDMA engines available for Write");

  info_.localMemSizePerCU_ = group_segment_size;
  info_.localMemSize_ = group_segment_size;

  info_.maxWorkItemDimensions_ = 3;

  uint8_t memory_properties[8];
  // Get the memory property from ROCr.
  if (HSA_STATUS_SUCCESS != hsa_agent_get_info(bkendDevice_,
                              (hsa_agent_info_t) HSA_AMD_AGENT_INFO_MEMORY_PROPERTIES,
                              memory_properties)) {
    LogError("HSA_AGENT_INFO_AMD_MEMORY_PROPERTIES query failed");
  }

  // Check if the device is APU
  if (hsa_flag_isset64(memory_properties, HSA_AMD_MEMORY_PROPERTY_AGENT_IS_APU)) {
    info_.hostUnifiedMemory_ = 1;
  }

  if (settings().enableLocalMemory_ && gpuvm_segment_.handle != 0) {
    size_t global_segment_size = 0;
    if (HSA_STATUS_SUCCESS != hsa_amd_memory_pool_get_info(gpuvm_segment_,
                                                           HSA_AMD_MEMORY_POOL_INFO_SIZE,
                                                           &global_segment_size)) {
      return false;
    }

    assert(global_segment_size > 0);
    info_.globalMemSize_ = (static_cast<uint64_t>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
                            static_cast<uint64_t>(global_segment_size)) / 100u;

    // For APU with vram size <= 512MiB, use a smaller single alloc percentage
    if (info_.globalMemSize_ <= 536870912) {
      if (flagIsDefault(GPU_SINGLE_ALLOC_PERCENT)) {
        GPU_SINGLE_ALLOC_PERCENT = 75;
      }
    }
    // Limit gpu single allocation percentage for gfx940
    if ((isa().versionMajor() == 9) && (isa().versionMinor() == 4) &&
        (isa().versionStepping() == 0) && (info_.hostUnifiedMemory_ == 1)) {
      if (gpu_agents_.size() == 1 || p2p_agents_.size() == 0) {
        if (flagIsDefault(GPU_SINGLE_ALLOC_PERCENT)) {
            GPU_SINGLE_ALLOC_PERCENT = 60;
        }
      }
    }

    gpuvm_segment_max_alloc_ =
        uint64_t(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u);
    assert(gpuvm_segment_max_alloc_ > 0);

    info_.maxMemAllocSize_ = static_cast<uint64_t>(gpuvm_segment_max_alloc_);

    if (HSA_STATUS_SUCCESS !=
        hsa_amd_memory_pool_get_info(gpuvm_segment_, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE,
                                     &alloc_granularity_)) {
      return false;
    }

    assert(alloc_granularity_ > 0);
  } else {
    // We suppose half of physical memory can be used by GPU in APU system
    info_.globalMemSize_ =
        uint64_t(sysconf(_SC_PAGESIZE)) * uint64_t(sysconf(_SC_PHYS_PAGES)) / 2;
    info_.globalMemSize_ = std::max(info_.globalMemSize_, uint64_t(1 * Gi));
    info_.globalMemSize_ = (static_cast<uint64_t>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
                            static_cast<uint64_t>(info_.globalMemSize_)) / 100u;

    info_.maxMemAllocSize_ =
        uint64_t(info_.globalMemSize_ * std::min(GPU_SINGLE_ALLOC_PERCENT, 100u) / 100u);

    if (HSA_STATUS_SUCCESS !=
        hsa_amd_memory_pool_get_info(
            system_segment_, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULE, &alloc_granularity_)) {
      return false;
    }
  }

  freeMem_ = info_.globalMemSize_;

  // Make sure the max allocation size is not larger than the available memory size.
  info_.maxMemAllocSize_ = std::min(info_.maxMemAllocSize_, info_.globalMemSize_);
  info_.maxMemAllocSize_ = amd::alignDown(info_.maxMemAllocSize_, sizeof(uint64_t));

  // Maximum system memory allocation size allowed
  info_.maxPhysicalMemAllocSize_ = amd::Os::getPhysicalMemSize();

  // make sure we don't run anything over 8 params for now
  info_.maxParameterSize_ = 1024;

  uint32_t max_work_group_size = 0;
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, &max_work_group_size)) {
    return false;
  }
  assert(max_work_group_size > 0);
  max_work_group_size =
      std::min(max_work_group_size, static_cast<uint32_t>(settings().maxWorkGroupSize_));
  info_.maxWorkGroupSize_ = max_work_group_size;

  uint16_t max_workgroup_size[3] = {0, 0, 0};
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_WORKGROUP_MAX_DIM, &max_workgroup_size)) {
    return false;
  }
  assert(max_workgroup_size[0] != 0 && max_workgroup_size[1] != 0 && max_workgroup_size[2] != 0);

  uint16_t max_work_item_size = static_cast<uint16_t>(max_work_group_size);
  info_.maxWorkItemSizes_[0] = std::min(max_workgroup_size[0], max_work_item_size);
  info_.maxWorkItemSizes_[1] = std::min(max_workgroup_size[1], max_work_item_size);
  info_.maxWorkItemSizes_[2] = std::min(max_workgroup_size[2], max_work_item_size);
  info_.preferredWorkGroupSize_ = settings().preferredWorkGroupSize_;

  info_.nativeVectorWidthChar_ = info_.preferredVectorWidthChar_ = 4;
  info_.nativeVectorWidthShort_ = info_.preferredVectorWidthShort_ = 2;
  info_.nativeVectorWidthInt_ = info_.preferredVectorWidthInt_ = 1;
  info_.nativeVectorWidthLong_ = info_.preferredVectorWidthLong_ = 1;
  info_.nativeVectorWidthFloat_ = info_.preferredVectorWidthFloat_ = 1;

  if (agent_profile_ == HSA_PROFILE_FULL) {  // full-profile = participating in coherent memory,
                                             // base-profile = NUMA based non-coherent memory
    info_.hostUnifiedMemory_ = 1;
    info_.iommuv2_ = true;
  }
  info_.memBaseAddrAlign_ = 8 * (flagIsDefault(MEMOBJ_BASE_ADDR_ALIGN) ?
      sizeof(int64_t[16]) * 2 : MEMOBJ_BASE_ADDR_ALIGN);
  info_.minDataTypeAlignSize_ = sizeof(int64_t[16]);

  info_.maxConstantArgs_ = 8;
  info_.preferredConstantBufferSize_ = 16 * Ki;
  info_.maxConstantBufferSize_ = info_.maxMemAllocSize_;
  info_.localMemType_ = CL_LOCAL;
  info_.errorCorrectionSupport_ = false;
  info_.profilingTimerResolution_ = 1;
  info_.littleEndian_ = true;
  info_.compilerAvailable_ = true;
  info_.executionCapabilities_ = CL_EXEC_KERNEL;
  info_.queueProperties_ = CL_QUEUE_PROFILING_ENABLE;
  info_.platform_ = AMD_PLATFORM;
  info_.profile_ = "FULL_PROFILE";
  ::strncpy(info_.vendor_, "Advanced Micro Devices, Inc.", sizeof(info_.vendor_) - 1);

  info_.addressBits_ = LP64_SWITCH(32, 64);
  info_.maxSamplers_ = 16;
  info_.bufferFromImageSupport_ = false;
  info_.oclcVersion_ = "OpenCL C " OPENCL_C_VERSION_STR " ";
  info_.spirVersions_ = "";

  uint16_t major, minor;
  if (hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_VERSION_MAJOR, &major) !=
          HSA_STATUS_SUCCESS ||
      hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_VERSION_MINOR, &minor) !=
          HSA_STATUS_SUCCESS) {
    return false;
  }
  std::stringstream ss;
  ss << AMD_BUILD_STRING " (HSA" << major << "." << minor << "," << (settings().useLightning_ ? "LC" : "HSAIL");
  ss <<  ")";

  ::strncpy(info_.driverVersion_, ss.str().c_str(), sizeof(info_.driverVersion_) - 1);

  if (isa().versionMajor() >= 9) {
    info_.version_ = "OpenCL " /*OPENCL_VERSION_STR*/"2.0" " ";
  } else {
    info_.version_ = "OpenCL " /*OPENCL_VERSION_STR*/"1.2" " ";
  }

  info_.builtInKernels_ = "";
  info_.linkerAvailable_ = true;
  info_.preferredInteropUserSync_ = true;
  info_.printfBufferSize_ = PrintfDbg::WorkitemDebugSize * info().maxWorkGroupSize_;
  info_.vendorId_ = 0x1002;  // AMD's PCIe vendor id

  info_.maxGlobalVariableSize_ = static_cast<size_t>(info_.maxMemAllocSize_);
  info_.globalVariablePreferredTotalSize_ = static_cast<size_t>(info_.globalMemSize_);

  // Populate the single config setting.
  info_.singleFPConfig_ =
      CL_FP_ROUND_TO_NEAREST | CL_FP_ROUND_TO_ZERO | CL_FP_ROUND_TO_INF | CL_FP_INF_NAN | CL_FP_FMA;

  if (settings().doublePrecision_) {
    info_.doubleFPConfig_ = info_.singleFPConfig_ | CL_FP_DENORM;
    info_.singleFPConfig_ |= CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT;
  }

  if (settings().singleFpDenorm_) {
    info_.singleFPConfig_ |= CL_FP_DENORM;
  }

  if (settings().checkExtension(ClKhrFp16)) {
    info_.halfFPConfig_ = info_.singleFPConfig_;
  }

  info_.preferredPlatformAtomicAlignment_ = 0;
  info_.preferredGlobalAtomicAlignment_ = 0;
  info_.preferredLocalAtomicAlignment_ = 0;

  uint8_t hsa_extensions[128];
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_EXTENSIONS, hsa_extensions)) {
    return false;
  }

  assert(HSA_EXTENSION_IMAGES < 8);
  const bool image_is_supported = ((hsa_extensions[0] & (1 << HSA_EXTENSION_IMAGES)) != 0);
  if (image_is_supported) {
    // Images
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_MAX_SAMPLER_HANDLERS),
                           &info_.maxSamplers_)) {
      return false;
    }

    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_MAX_IMAGE_RD_HANDLES),
                           &info_.maxReadImageArgs_)) {
      return false;
    }

    // TODO: no attribute for write image.
    info_.maxWriteImageArgs_ = 8;

    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_MAX_IMAGE_RORW_HANDLES),
                           &info_.maxReadWriteImageArgs_)) {
      return false;
    }

    uint32_t image_max_dim[3];
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_IMAGE_2D_MAX_ELEMENTS),
                           &image_max_dim)) {
      return false;
    }

    info_.image2DMaxWidth_ = image_max_dim[0];
    info_.image2DMaxHeight_ = image_max_dim[1];

    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS),
                           &image_max_dim)) {
      return false;
    }

    info_.image3DMaxWidth_ = image_max_dim[0];
    info_.image3DMaxHeight_ = image_max_dim[1];
    info_.image3DMaxDepth_ = image_max_dim[2];

    uint32_t max_array_size = 0;
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_IMAGE_ARRAY_MAX_LAYERS),
                           &max_array_size)) {
      return false;
    }

    info_.imageMaxArraySize_ = max_array_size;

    uint32_t max_image1da_width = 0;
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_IMAGE_1DA_MAX_ELEMENTS),
                           &max_image1da_width)) {
      return false;
    }

    info_.image1DAMaxWidth_ = max_image1da_width;

    uint32_t max_image2da_width[2] = {0, 0};
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_IMAGE_2DA_MAX_ELEMENTS),
                           &max_image2da_width)) {
      return false;
    }

    info_.image2DAMaxWidth_[0] = max_image2da_width[0];
    info_.image2DAMaxWidth_[1] = max_image2da_width[1];

    uint32_t max_image1d_width = 0;
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS),
                           &max_image1d_width)) {
      return false;
    }
    info_.image1DMaxWidth_ = max_image1d_width;

    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_EXT_AGENT_INFO_IMAGE_1DB_MAX_ELEMENTS),
                           &image_max_dim)) {
      return false;
    }
    info_.imageMaxBufferSize_ = (amd::IS_HIP) ? image_max_dim[0] : (1 << 27);

    info_.imagePitchAlignment_ = 256;

    info_.imageBaseAddressAlignment_ = 256;

    info_.bufferFromImageSupport_ = false;

    info_.imageSupport_ = (info_.maxReadWriteImageArgs_ > 0) ? true : false;
  }

  // Enable SVM Capabilities of Hsa device. Ensure
  // user has not setup memory to be non-coherent
  info_.svmCapabilities_ = 0;
  if (!settings().enableNCMode_) {
    info_.svmCapabilities_ = CL_DEVICE_SVM_COARSE_GRAIN_BUFFER;
    info_.svmCapabilities_ |= CL_DEVICE_SVM_FINE_GRAIN_BUFFER;
    // Report fine-grain system only on full profile
    if (agent_profile_ == HSA_PROFILE_FULL) {
      info_.svmCapabilities_ |= CL_DEVICE_SVM_FINE_GRAIN_SYSTEM;
    }
    if (amd::IS_HIP) {
      if (info_.iommuv2_ || isa().versionMajor() >= 8) {
        info_.svmCapabilities_ |= CL_DEVICE_SVM_ATOMICS;
      }
    }
    else if (!settings().useLightning_) {
      if (info_.iommuv2_ || (isa().versionMajor() == 8)) {
        info_.svmCapabilities_ |= CL_DEVICE_SVM_ATOMICS;
      }
    }
  }

  if (settings().checkExtension(ClAmdDeviceAttributeQuery)) {
    info_.simdWidth_ = isa().simdWidth();
    info_.simdInstructionWidth_ = isa().simdInstructionWidth();
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_WAVEFRONT_SIZE, &info_.wavefrontWidth_)) {
      return false;
    }

    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MEMORY_WIDTH),
                           &info_.vramBusBitWidth_)) {
      return false;
    }

    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_NUM_SIMDS_PER_CU),
                           &info_.simdPerCU_)) {
      return false;
    }

    uint32_t max_waves_per_cu = 0;
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU),
                           &max_waves_per_cu)) {
      return false;
    }

    if (settings().enableWgpMode_) {
      info_.simdPerCU_ *= 2;
      max_waves_per_cu *= 2;
    }

    info_.maxThreadsPerCU_ = info_.wavefrontWidth_ * max_waves_per_cu;
    uint32_t cache_sizes[4];
    /* FIXIT [skudchad] -  Seems like hardcoded in HSA backend so 0*/
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_AGENT_INFO_CACHE_SIZE),
                           cache_sizes)) {
      return false;
    }

    uint32_t asic_revision = 0;
    if (HSA_STATUS_SUCCESS !=
        hsa_agent_get_info(bkendDevice_,
                           static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_ASIC_REVISION),
                           &asic_revision)) {
      return false;
    }
    info_.asicRevision_ = asic_revision;

    info_.l2CacheSize_ = cache_sizes[1];
    info_.timeStampFrequency_ = 1000000;
    info_.globalMemChannelBanks_ = 4;
    info_.globalMemChannelBankWidth_ = isa().memChannelBankWidth();
    info_.localMemSizePerCU_ = isa().localMemSizePerCU();
    info_.localMemBanks_ = isa().localMemBanks();
    info_.numAsyncQueues_ = kMaxAsyncQueues;
    info_.numRTQueues_ = info_.numAsyncQueues_;
    info_.numRTCUs_ = info_.maxComputeUnits_;

    //TODO: set to true once thread trace support is available
    info_.threadTraceEnable_ = false;
    info_.pcieDeviceId_ = pciDeviceId_;
    info_.cooperativeGroups_ = settings().enableCoopGroups_;
    info_.cooperativeMultiDeviceGroups_ = settings().enableCoopMultiDeviceGroups_;
    // Enable StreamWrite and StreamWait for all devices
    info_.aqlBarrierValue_ = true;
  }

  info_.maxPipePacketSize_ = info_.maxMemAllocSize_;
  info_.maxPipeActiveReservations_ = 16;
  info_.maxPipeArgs_ = 16;

  info_.queueOnDeviceProperties_ =
      CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE;
  info_.queueOnDevicePreferredSize_ = 256 * Ki;
  info_.queueOnDeviceMaxSize_ = 8 * Mi;
  info_.maxOnDeviceQueues_ = 1;
  info_.maxOnDeviceEvents_ = settings().numDeviceEvents_;

  // Get Values from from Comgr
  amd_comgr_metadata_node_t isaMeta;
  if (getIsaMeta(std::move(isa().isaName()), isaMeta)) {
    std::string vgprValue;
    info_.availableVGPRs_ = (getValueFromIsaMeta(isaMeta, "AddressableNumVGPRs", vgprValue))
        ? (atoi(vgprValue.c_str()) * info_.simdPerCU_)
        : 0;

    info_.availableRegistersPerCU_ = info_.availableVGPRs_ * 64;  // 64 registers per VGPR

    std::string sgprValue;
    info_.availableSGPRs_ = (getValueFromIsaMeta(isaMeta, "AddressableNumSGPRs", sgprValue))
        ? (atoi(sgprValue.c_str()))
        : 0;
    if (!releaseIsaMeta(isaMeta)) {
      LogInfo("Can not release the isa meta node");
    }
  }

  // Generic support for HMM interfaces
  if (HSA_STATUS_SUCCESS != hsa_system_get_info(HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED,
      &info_.hmmSupported_)) {
    LogError("HSA_AMD_SYSTEM_INFO_SVM_SUPPORTED query failed. HMM will be disabled");
  }

  // This capability should be available with xnack enabled
  if (HSA_STATUS_SUCCESS != hsa_system_get_info(HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT,
      &info_.hmmCpuMemoryAccessible_)) {
    LogError("HSA_AMD_SYSTEM_INFO_SVM_ACCESSIBLE_BY_DEFAULT query failed.");
  }

  // HMM specific capability for CPU direct access to device memory
  if (HSA_STATUS_SUCCESS != hsa_agent_get_info(bkendDevice_,
      static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS),
      &info_.hmmDirectHostAccess_)) {
    LogError("HSA_AMD_AGENT_INFO_SVM_DIRECT_HOST_ACCESS query failed.");
  }

  ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Gfx Major/Minor/Stepping: %d/%d/%d", isa().versionMajor(),
  isa().versionMinor(), isa().versionStepping());
  ClPrint(amd::LOG_INFO, amd::LOG_INIT, "HMM support: %d, XNACK: %d, Direct host access: %d",
    info_.hmmSupported_, info_.hmmCpuMemoryAccessible_, info_.hmmDirectHostAccess_);
  ClPrint(amd::LOG_INFO, amd::LOG_INIT, "Max SDMA Read Mask: 0x%x, Max SDMA Write Mask: 0x%x",
          maxSdmaReadMask_, maxSdmaWriteMask_);

  info_.globalCUMask_ = {};

  // Virtual memory Management Support, if set to true then the HW and SW Stack supports VMM.
  info_.virtualMemoryManagement_ = false;
  if (HIP_VMEM_MANAGE_SUPPORT) {
    if (HSA_STATUS_SUCCESS != hsa_system_get_info(
          static_cast<hsa_system_info_t>(HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED),
          &info_.virtualMemoryManagement_)) {
      LogError("HSA_AMD_SYSTEM_INFO_VIRTUAL_MEM_API_SUPPORTED query failed ");
    }
  }
  HIP_MEM_POOL_USE_VM &= info_.virtualMemoryManagement_;

  switch (isa().versionMajor()) {
    case (11):
      if (isa().versionMinor() == 0) {
        switch (isa().versionStepping()) {
          case (0):
          case (1):
            info_.vgprAllocGranularity_ = 24;
            info_.vgprsPerSimd_ = 1536;
            break;
          case (2):
          case (3):
          default:
            info_.vgprAllocGranularity_ = 16;
            info_.vgprsPerSimd_ = 1024;
            break;
        }
      } else if (isa().versionMinor() == 5) {
        switch (isa().versionStepping()) {
          case (1):
            info_.vgprAllocGranularity_ = 24;
            info_.vgprsPerSimd_ = 1536;
            break;
          default:
            info_.vgprAllocGranularity_ = 16;
            info_.vgprsPerSimd_ = 1024;
            break;
        }
      } else {
          info_.vgprAllocGranularity_ = 16;
          info_.vgprsPerSimd_ = 1024;
      }
      break;
    case (10):
      switch (isa().versionMinor()) {
        case (0):
        case (1):
          info_.vgprAllocGranularity_ = 8;
          info_.vgprsPerSimd_ = 1024;
          break;
        case (3):
        default:
          info_.vgprAllocGranularity_ = 16;
          info_.vgprsPerSimd_ = 1024;
          break;
      }
      break;
    case (9):
      if ((isa().versionMinor() == 0 && isa().versionStepping() == 10) ||
          (isa().versionMinor() == 4)) {
        info_.vgprAllocGranularity_ = 8;
        info_.vgprsPerSimd_ = 512;
      } else {
        info_.vgprAllocGranularity_ = 4;
        info_.vgprsPerSimd_ = 256;
      }
      break;
    default:
      // For gfx<=8
      info_.vgprAllocGranularity_ = 4;
      info_.vgprsPerSimd_ = 256;
      break;
  }

  if (isa().versionMajor() < 8) {
    info_.sgprsPerSimd_ = 512;
  } else if (isa().versionMajor() < 10) {
    info_.sgprsPerSimd_ = 800;
  } else {
    info_.sgprsPerSimd_ =
        std::numeric_limits<uint32_t>::max();  // gfx10+ does not share SGPRs between waves
  }

  return true;
}

// ================================================================================================
device::VirtualDevice* Device::createVirtualDevice(amd::CommandQueue* queue) {
  amd::ScopedLock lock(vgpusAccess());

  bool profiling = (queue != nullptr) && queue->properties().test(CL_QUEUE_PROFILING_ENABLE);
  bool cooperative = false;

  // If amd command queue is null, then it's an internal device queue
  if (queue == nullptr) {
    // In HIP mode the device queue will be allocated for the cooperative launches only
    cooperative = amd::IS_HIP && settings().enableCoopGroups_;
    profiling = amd::IS_HIP;
  }
  // Initialization of heap and other resources occur during the command
  // queue creation time.
  const std::vector<uint32_t> defaultCuMask = {};
  bool q = (queue != nullptr);
  VirtualGPU* virtualDevice = new VirtualGPU(*this, profiling, cooperative,
                                            q ? queue->cuMask() : defaultCuMask,
                                            q ? queue->priority()
                                              : amd::CommandQueue::Priority::Normal);

  if (!virtualDevice->create()) {
    delete virtualDevice;
    return nullptr;
  }

  return virtualDevice;
}

bool Device::globalFreeMemory(size_t* freeMemory) const {
  const uint TotalFreeMemory = 0;
  const uint LargestFreeBlock = 1;
  uint64_t globalAvailMemory;
  // Queries memory available in bytes across all global pools owned by the agent
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_,
                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MEMORY_AVAIL),
                         &globalAvailMemory)) {
    LogError("HSA_AMD_AGENT_INFO_MEMORY_AVAIL query failed.");
    return false;
  }

  globalAvailMemory = globalAvailMemory / Ki;
  if (globalAvailMemory > HIP_HIDDEN_FREE_MEM * Ki) {
    globalAvailMemory -= HIP_HIDDEN_FREE_MEM * Ki;
  } else {
    globalAvailMemory = 0;
  }

  freeMemory[TotalFreeMemory] = globalAvailMemory;
  // since there is no memory heap on ROCm, the biggest free block is
  // equal to total free local memory
  freeMemory[LargestFreeBlock] = freeMemory[TotalFreeMemory];

  return true;
}

bool Device::bindExternalDevice(uint flags, void* const gfxDevice[], void* gfxContext,
                                bool validateOnly) {
#if defined(_WIN32)
  return false;
#else
  if ((flags & amd::Context::GLDeviceKhr) == 0) return false;

  MesaInterop::MESA_INTEROP_KIND kind = MesaInterop::MESA_INTEROP_NONE;
  MesaInterop::DisplayHandle display;
  MesaInterop::ContextHandle context;

  if ((flags & amd::Context::EGLDeviceKhr) != 0) {
    kind = MesaInterop::MESA_INTEROP_EGL;
    display.eglDisplay = reinterpret_cast<EGLDisplay>(gfxDevice[amd::Context::GLDeviceKhrIdx]);
    context.eglContext = reinterpret_cast<EGLContext>(gfxContext);
  } else {
    kind = MesaInterop::MESA_INTEROP_GLX;
    display.glxDisplay = reinterpret_cast<Display*>(gfxDevice[amd::Context::GLDeviceKhrIdx]);
    context.glxContext = reinterpret_cast<GLXContext>(gfxContext);
  }

  mesa_glinterop_device_info info;
  info.version = MESA_GLINTEROP_DEVICE_INFO_VERSION;
  if (!MesaInterop::Init(kind)) {
    return false;
  }

  if (!MesaInterop::GetInfo(info, kind, display, context)) {
    return false;
  }

  return info_.deviceTopology_.pcie.bus == info.pci_bus &&
      info_.deviceTopology_.pcie.device == info.pci_device &&
      info_.deviceTopology_.pcie.function == info.pci_function &&
      info_.vendorId_ == info.vendor_id && pciDeviceId_ == info.device_id;

#endif
}

bool Device::unbindExternalDevice(uint flags, void* const gfxDevice[], void* gfxContext,
                                  bool validateOnly) {
#if defined(_WIN32)
  return false;
#else
  if ((flags & amd::Context::GLDeviceKhr) == 0) return false;
  return true;
#endif
}

amd::Memory* Device::findMapTarget(size_t size) const {
  // Must be serialised for access
  amd::ScopedLock lk(*mapCacheOps_);

  amd::Memory* map = nullptr;
  size_t minSize = 0;
  size_t maxSize = 0;
  uint mapId = mapCache_->size();
  uint releaseId = mapCache_->size();

  // Find if the list has a map target of appropriate size
  for (uint i = 0; i < mapCache_->size(); i++) {
    if ((*mapCache_)[i] != nullptr) {
      // Requested size is smaller than the entry size
      if (size < (*mapCache_)[i]->getSize()) {
        if ((minSize == 0) || (minSize > (*mapCache_)[i]->getSize())) {
          minSize = (*mapCache_)[i]->getSize();
          mapId = i;
        }
      }
      // Requeted size matches the entry size
      else if (size == (*mapCache_)[i]->getSize()) {
        mapId = i;
        break;
      } else {
        // Find the biggest map target in the list
        if (maxSize < (*mapCache_)[i]->getSize()) {
          maxSize = (*mapCache_)[i]->getSize();
          releaseId = i;
        }
      }
    }
  }

  // Check if we found any map target
  if (mapId < mapCache_->size()) {
    map = (*mapCache_)[mapId];
    (*mapCache_)[mapId] = nullptr;
  }
  // If cache is full, then release the biggest map target
  else if (releaseId < mapCache_->size()) {
    (*mapCache_)[releaseId]->release();
    (*mapCache_)[releaseId] = nullptr;
  }

  return map;
}

bool Device::addMapTarget(amd::Memory* memory) const {
  // Must be serialised for access
  amd::ScopedLock lk(*mapCacheOps_);

  // the svm memory shouldn't be cached
  if (!memory->canBeCached()) {
    return false;
  }
  // Find if the list has a map target of appropriate size
  for (uint i = 0; i < mapCache_->size(); ++i) {
    if ((*mapCache_)[i] == nullptr) {
      (*mapCache_)[i] = memory;
      return true;
    }
  }

  // Add a new entry
  mapCache_->push_back(memory);

  return true;
}

Memory* Device::getRocMemory(amd::Memory* mem) const {
  return static_cast<roc::Memory*>(mem->getDeviceMemory(*this));
}

// ================================================================================================
device::Memory* Device::createMemory(amd::Memory& owner) const {
  roc::Memory* memory = nullptr;
  if (owner.asBuffer()) {
    memory = new roc::Buffer(*this, owner);
  } else if (owner.asImage()) {
    memory = new roc::Image(*this, owner);
  } else {
    LogError("Unknown memory type");
  }

  if (memory == nullptr) {
    return nullptr;
  }

  bool result = memory->create();

  if (!result) {
    LogError("Failed creating memory");
    delete memory;
    return nullptr;
  }

  if (isP2pEnabled()) {
    memory->setAllowedPeerAccess(true);
  }
  // Initialize if the memory is a pipe object
  if (owner.getType() == CL_MEM_OBJECT_PIPE) {
    // Pipe initialize in order read_idx, write_idx, end_idx. Refer clk_pipe_t structure.
    // Init with 3 DWORDS for 32bit addressing and 6 DWORDS for 64bit
    size_t pipeInit[3] = { 0, 0, owner.asPipe()->getMaxNumPackets() };
    xferMgr().writeBuffer(pipeInit, *memory, amd::Coord3D(0), amd::Coord3D(sizeof(pipeInit)));
  }

  // Transfer data only if OCL context has one device.
  // Cache coherency layer will update data for multiple devices
  if (!memory->isHostMemDirectAccess() && owner.asImage() && (owner.parent() == nullptr) &&
      (owner.getMemFlags() & CL_MEM_COPY_HOST_PTR) &&
      (owner.getContext().devices().size() == 1)) {
    // To avoid recurssive call to Device::createMemory, we perform
    // data transfer to the view of the image
    amd::Image* imageView = owner.asImage()->createView(owner.getContext(),
        owner.asImage()->getImageFormat(), xferQueue());

    if (imageView == nullptr) {
      LogError("[OCL] Fail to allocate view of image object");
      return nullptr;
    }

    Image* devImageView = new roc::Image(static_cast<const Device&>(*this), *imageView);
    if (devImageView == nullptr) {
      LogError("[OCL] Fail to allocate device mem object for the view");
      imageView->release();
      return nullptr;
    }

    if (devImageView != nullptr && !devImageView->createView(static_cast<roc::Image&>(*memory))) {
      LogError("[OCL] Fail to create device mem object for the view");
      delete devImageView;
      imageView->release();
      return nullptr;
    }

    imageView->replaceDeviceMemory(this, devImageView);

    // Copy data with the original pitch values, since runtime doesn't perform
    // extra sysmem allocation for one device
    const auto image = owner.asImage();
    result = xferMgr().writeImage(owner.getHostMem(), *devImageView, amd::Coord3D(0, 0, 0),
        imageView->getRegion(), image->getRowPitch(), image->getSlicePitch(), true);

    // Release host memory, since runtime copied data
    owner.setHostMem(nullptr);

    imageView->release();
  }

  // Prepin sysmem buffer for possible data synchronization between CPU and GPU
  if (!memory->isHostMemDirectAccess() &&
      // Pin memory for the parent object only
      (owner.parent() == nullptr) &&
      (owner.getHostMem() != nullptr) &&
      (owner.getSvmPtr() == nullptr)) {
    memory->pinSystemMemory(owner.getHostMem(), owner.getSize());
  }

  if (!result) {
    delete memory;
    DevLogError("Cannot Write Image \n");
    return nullptr;
  }

  return memory;
}

// ================================================================================================
device::Memory* Device::createMemory(size_t size) const {
  auto buffer = new roc::Buffer(*this, size);
  static constexpr bool LocalAlloc = true;
  if ((buffer == nullptr) || !buffer->create(LocalAlloc)) {
    LogError("Couldn't allocate memory on device!");
    return nullptr;
  }
  return buffer;
}

// ================================================================================================
void* Device::hostAlloc(size_t size, size_t alignment, MemorySegment mem_seg) const {
  void* ptr = nullptr;

  hsa_amd_memory_pool_t segment{0};
  switch (mem_seg) {
    case kKernArg : {
      if (settings().fgs_kernel_arg_) {
        segment = system_kernarg_segment_;
        break;
      }
      // Falls through on else case.
    }
    case kNoAtomics :
      // If runtime disables barrier, then all host allocations must have L2 disabled
      if (system_coarse_segment_.handle != 0) {
        segment = system_coarse_segment_;
        break;
      }
      // Falls through on else case.
    case kAtomics :
      segment = system_segment_;
      break;
    default :
      guarantee(false, "Invalid Memory Segment");
      break;
  }

  assert(segment.handle != 0);
  hsa_status_t stat = hsa_amd_memory_pool_allocate(segment, size, 0, &ptr);
  ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Allocate hsa host memory %p, size 0x%zx,"
          " numa_node = %d", ptr, size, preferred_numa_node_);
  if (stat != HSA_STATUS_SUCCESS) {
    LogPrintfError("Fail allocation host memory with err %d", stat);
    return nullptr;
  }

  stat = hsa_amd_agents_allow_access(gpu_agents_.size(), &gpu_agents_[0], nullptr, ptr);
  if (stat != HSA_STATUS_SUCCESS) {
    LogPrintfError("Fail hsa_amd_agents_alloc_access with err %d", stat);
    hostFree(ptr, size);
    return nullptr;
  }

  return ptr;
}

// ================================================================================================
void* Device::hostAgentAlloc(size_t size, const AgentInfo& agentInfo, bool atomics) const {
  void* ptr = nullptr;
  const hsa_amd_memory_pool_t segment =
      // If runtime disables barrier, then all host allocations must have L2 disabled
      !atomics ? (agentInfo.coarse_grain_pool.handle != 0) ?
              agentInfo.coarse_grain_pool : agentInfo.fine_grain_pool
               : agentInfo.fine_grain_pool;
  assert(segment.handle != 0);
  hsa_status_t stat = hsa_amd_memory_pool_allocate(segment, size, 0, &ptr);
  ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Allocate hsa host memory %p, size 0x%zx", ptr, size);
  if (stat != HSA_STATUS_SUCCESS) {
    LogPrintfError("Fail allocation host memory with err %d", stat);
    return nullptr;
  }

  stat = hsa_amd_agents_allow_access(gpu_agents_.size(), &gpu_agents_[0], nullptr, ptr);
  if (stat != HSA_STATUS_SUCCESS) {
    LogPrintfError("Fail hsa_amd_agents_allow_access with err %d", stat);
    hostFree(ptr, size);
    return nullptr;
  }

  return ptr;
}

// ================================================================================================
void* Device::hostNumaAlloc(size_t size, size_t alignment, bool atomics) const {
  void* ptr = nullptr;
#ifndef ROCCLR_SUPPORT_NUMA_POLICY
  ptr = hostAlloc(size, alignment, atomics
                  ? Device::MemorySegment::kAtomics : Device::MemorySegment::kNoAtomics);
#else
  int mode = MPOL_DEFAULT;
  int maxNodes = numa_num_possible_nodes();
  bitmask* nodeMask = numa_bitmask_alloc(maxNodes);
  auto cpuCount = cpu_agents_.size();

  long res = get_mempolicy(&mode, nodeMask->maskp, nodeMask->size, NULL, 0);
  if (res) {
    LogPrintfError("get_mempolicy failed with error %ld", res);
    return ptr;
  }
  ClPrint(amd::LOG_INFO, amd::LOG_RESOURCE,
          "get_mempolicy() succeed with mode %d, nodeMask 0x%lx, cpuCount %zu",
          mode, *nodeMask->maskp, cpuCount);

  switch (mode) {
    // For details, see "man get_mempolicy".
    case MPOL_BIND:
    case MPOL_PREFERRED:
      // We only care about the first CPU node
      for (unsigned int i = 0; i < cpuCount; i++) {
        if ((1u << i) & *nodeMask->maskp) {
          ptr = hostAgentAlloc(size, cpu_agents_[i], atomics);
          break;
        }
      }
      break;
    default:
      //  All other modes fall back to default mode
      ptr = hostAlloc(size, alignment, atomics
                      ? Device::MemorySegment::kAtomics : Device::MemorySegment::kNoAtomics);
  }
  numa_free_cpumask(nodeMask);
#endif // ROCCLR_SUPPORT_NUMA_POLICY
  return ptr;
}

void Device::hostFree(void* ptr, size_t size) const { memFree(ptr, size); }

bool Device::deviceAllowAccess(void* ptr) const {
  std::lock_guard<std::mutex> lock(lock_allow_access_);
  if (!p2pAgents().empty()) {
    hsa_status_t stat = hsa_amd_agents_allow_access(p2pAgents().size(),
                                                    p2pAgents().data(), nullptr, ptr);
    if (stat != HSA_STATUS_SUCCESS) {
      LogPrintfError("Allow p2p access failed - hsa_amd_agents_allow_access with err %d", stat);
      return false;
    }
  }
  return true;
}

bool Device::allowPeerAccess(device::Memory* memory) const {
  if (memory == nullptr) {
    return false;
  }
  if (!p2pAgents().empty()) {
    void* ptr = reinterpret_cast<void*>(memory->virtualAddress());
    hsa_agent_t agent = getBackendDevice();
    hsa_status_t stat = hsa_amd_agents_allow_access(1, &agent, nullptr, ptr);
    if (stat != HSA_STATUS_SUCCESS) {
      LogPrintfError("Allow p2p access failed - hsa_amd_agents_allow_access with err: %d", stat);
      return false;
    }
  }
  return true;
}

uint64_t Device::deviceVmemAlloc(size_t size, uint64_t flags) const {
  hsa_amd_vmem_alloc_handle_t hsa_vmem_handle {};

  // We only allow pinned memory at this time.
  hsa_status_t hsa_status = hsa_amd_vmem_handle_create(gpuvm_segment_, size, MEMORY_TYPE_PINNED,
                                                       flags, &hsa_vmem_handle);
  if (hsa_status != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_amd_vmem_handle_create! Failed with hsa status: %d \n", hsa_status);
  }

  return hsa_vmem_handle.handle;
}

void Device::deviceVmemRelease(uint64_t mem_handle) const {
  hsa_amd_vmem_alloc_handle_t hsa_vmem_handle {};
  hsa_vmem_handle.handle = mem_handle;

  hsa_status_t hsa_status = hsa_amd_vmem_handle_release(hsa_vmem_handle);
  if (hsa_status != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_amd_vmem_handle_release! Failed with hsa status: %d \n", hsa_status);
  }
}

void* Device::deviceLocalAlloc(size_t size, bool atomics, bool pseudo_fine_grain,
                               bool contiguous) const {

  const hsa_amd_memory_pool_t& pool = (pseudo_fine_grain && gpu_ext_fine_grained_segment_.handle)
                                       ? gpu_ext_fine_grained_segment_
                                         : (atomics && gpu_fine_grained_segment_.handle)
                                            ? gpu_fine_grained_segment_ : gpuvm_segment_;

  if (pool.handle == 0 || gpuvm_segment_max_alloc_ == 0) {
    DevLogPrintfError("Invalid argument, pool_handle: 0x%x , max_alloc: %u \n",
                      pool.handle, gpuvm_segment_max_alloc_);
    return nullptr;
  }

  uint32_t hsa_mem_flags = 0;
  if (contiguous) {
    hsa_mem_flags = HSA_AMD_MEMORY_POOL_CONTIGUOUS_FLAG;
  }

  void* ptr = nullptr;
  hsa_status_t stat = hsa_amd_memory_pool_allocate(pool, size, 0, &ptr);
  ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Allocate hsa device memory %p, size 0x%zx", ptr, size);
  if (stat != HSA_STATUS_SUCCESS) {
    LogError("Fail allocation local memory");
    return nullptr;
  }

  if (isP2pEnabled() && deviceAllowAccess(ptr) == false) {
    LogError("Allow p2p access for memory allocation");
    memFree(ptr, size);
    return nullptr;
  }
  return ptr;
}

void Device::memFree(void* ptr, size_t size) const {
  hsa_status_t stat = hsa_amd_memory_pool_free(ptr);
  ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, "Free hsa memory %p", ptr);
  if (stat != HSA_STATUS_SUCCESS) {
    LogError("Fail freeing local memory");
  }
}

void Device::updateFreeMemory(size_t size, bool free) {
  if (free) {
    freeMem_ += size;
  }
  else {
    if (size > freeMem_) {
      // To avoid underflow of the freeMem_
      // This can happen if the free mem tracked is inaccurate, as some allocations can happen
      // directly via ROCr
      ClPrint(amd::LOG_ERROR, amd::LOG_ALWAYS,
             "Free memory set to zero on device 0x%lx, requested size = 0x%zx, freeMem_ = 0x%zx",
             this, size, freeMem_.load());
      freeMem_ = 0;
      return;
    }
    freeMem_ -= size;
  }
  ClPrint(amd::LOG_INFO, amd::LOG_MEM, "Device=0x%lx, freeMem_ = 0x%zx", this, freeMem_.load());
}

// ================================================================================================
void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_mem_flags flags,
                       void* svmPtr) const {
  amd::Memory* mem = nullptr;

  if (nullptr == svmPtr) {
    // create a hidden buffer, which will allocated on the device later
    mem = new (context) amd::Buffer(context, flags, size,
                reinterpret_cast<void*>(amd::Memory::MemoryType::kSvmMemoryPtr));
    if (mem == nullptr) {
      LogError("failed to create a svm mem object!");
      return nullptr;
    }

    if (!mem->create(nullptr)) {
      LogError("failed to create a svm hidden buffer!");
      mem->release();
      return nullptr;
    }
    // if the device supports SVM FGS, return the committed CPU address directly.
    Memory* gpuMem = getRocMemory(mem);
    if (gpuMem == nullptr) {
      LogError("failed to create GPU memory from svm hidden buffer!");
      return nullptr;
    }

    // add the information to context so that we can use it later.
    if (mem->getSvmPtr() != nullptr) {
      amd::MemObjMap::AddMemObj(mem->getSvmPtr(), mem);
    }
    svmPtr = mem->getSvmPtr();
  } else {
    // Find the existing amd::mem object
    mem = amd::MemObjMap::FindMemObj(svmPtr);
    if (nullptr == mem) {
      DevLogPrintfError("Cannot find svm_ptr: 0x%x \n", svmPtr);
      return nullptr;
    }

    svmPtr = mem->getSvmPtr();
  }

  return svmPtr;
}

void* Device::virtualAlloc(void* req_addr, size_t size, size_t alignment) {
  void* vptr = nullptr;
  // Reserves the address using HSA APIs, with requested address.
  // There is no guarantee that we will get the requested address.
  hsa_status_t hsa_status = hsa_amd_vmem_address_reserve(&vptr, size,
                              reinterpret_cast<uint64_t>(req_addr), 0);
  if (hsa_status != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_amd_vmem_address_reserve. Failed with status: %d \n", hsa_status);
    return nullptr;
  }

  constexpr bool kParent = true;
  amd::Memory* mem = CreateVirtualBuffer(context(), vptr, size, -1, kParent);
  if (mem == nullptr) {
    LogPrintfError("Cannot create Virtual Buffer for vptr: %p of size: %u", vptr, size);
  }

  return mem->getSvmPtr();
}

void Device::virtualFree(void* addr) {
  amd::Memory* memObj = amd::MemObjMap::FindVirtualMemObj(addr);
  if (memObj == nullptr) {
    LogPrintfError("Cannot find the Virtual MemObj entry for this addr 0x%x", addr);
  }

  memObj->getContext().devices()[0]->DestroyVirtualBuffer(memObj);

  hsa_status_t hsa_status = hsa_amd_vmem_address_free(memObj->getSvmPtr(), memObj->getSize());
  if (hsa_status != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_amd_vmem_address_free. Failed with status:%d \n", hsa_status);
  }
}

bool Device::SetMemAccess(void* va_addr, size_t va_size, VmmAccess access_flags) {
  hsa_status_t hsa_status = HSA_STATUS_SUCCESS;
  hsa_amd_memory_access_desc_t desc;
  desc.permissions = static_cast<hsa_access_permission_t>(access_flags);
  desc.agent_handle = getBackendDevice();

  if ((hsa_status = hsa_amd_vmem_set_access(va_addr, va_size, &desc, 1))
                                                      != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_amd_vmem_set_access. Failed with status:%d \n", hsa_status);
    return false;
  }

  return true;
}

bool Device::GetMemAccess(void* va_addr, VmmAccess* access_flags_ptr) {
  hsa_status_t hsa_status = HSA_STATUS_SUCCESS;
  hsa_access_permission_t perms;

  size_t discard_offset = 0;
  amd::Memory* va_mem_obj = amd::MemObjMap::FindMemObj(va_addr, &discard_offset);
  if (va_mem_obj == nullptr) {
    LogPrintfError("Failed to get Memory Object for va_addr: 0x%x", va_addr);
    return false;
  }

  if ((hsa_status = hsa_amd_vmem_get_access(va_mem_obj->getSvmPtr(), &perms, getBackendDevice()))
                    != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_amd_vmem_get_access. Failed with status:%d \n", hsa_status);
    return false;
  }

  *access_flags_ptr = static_cast<VmmAccess>(perms);

  return true;
}

// ================================================================================================
bool Device::ExportShareableVMMHandle(amd::Memory& amd_mem_obj, int flags, void* shareableHandle) {
  hsa_status_t hsa_status = HSA_STATUS_SUCCESS;
  hsa_amd_vmem_alloc_handle_t hsa_vmem_handle {};
  hsa_vmem_handle.handle = amd_mem_obj.getUserData().hsa_handle;
  int dmabuf_fd = 0;

  if (hsa_vmem_handle.handle == 0) {
    LogError("HSA Handle is not valid");
    return false;
  }

  if ((hsa_status = hsa_amd_vmem_export_shareable_handle(&dmabuf_fd,
                      hsa_vmem_handle, flags)) != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_vmem_export_shareable_handle with status: %d \n", hsa_status);
    return false;
  }

  *(reinterpret_cast<int*>(shareableHandle)) = dmabuf_fd;

  return true;
}

// ================================================================================================
bool Device::ImportShareableHSAHandle(void* osHandle, uint64_t* hsa_handle_ptr) const {
  hsa_status_t hsa_status = HSA_STATUS_SUCCESS;
  hsa_amd_vmem_alloc_handle_t hsa_vmem_handle {};

  if (hsa_handle_ptr == nullptr) {
    LogError("HSA Handle ptr is null");
    return false;
  }

  int dmabuf_fd = *(reinterpret_cast<int*>(osHandle));
  if ((hsa_status = hsa_amd_vmem_import_shareable_handle(dmabuf_fd, &hsa_vmem_handle))
                      != HSA_STATUS_SUCCESS) {
    LogPrintfError("Failed hsa_amd_vmem_import_shareable_handle with status: %d \n", hsa_status);
    return false;
  }

  *hsa_handle_ptr = hsa_vmem_handle.handle;
  return true;
}

// ================================================================================================
amd::Memory* Device::ImportShareableVMMHandle(void* osHandle) {
  amd::Memory* amd_mem_obj = new (context()) amd::Buffer(context(),
                                  ROCCLR_MEM_PHYMEM | ROCCLR_MEM_INTERPROCESS, 0, osHandle);
  if (amd_mem_obj == nullptr) {
    LogError("Cannot create memory object");
    return nullptr;
  }

  if (!amd_mem_obj->create(nullptr, false)) {
    LogError("Failed to create mem_obj from imported fd");
    amd_mem_obj->release();
    return nullptr;
  }

  return amd_mem_obj;
}

// ================================================================================================
bool Device::SetSvmAttributesInt(const void* dev_ptr, size_t count,
                              amd::MemoryAdvice advice, bool first_alloc, bool use_cpu) const {
  if ((settings().hmmFlags_ & Settings::Hmm::EnableSvmTracking) && !first_alloc) {
    amd::Memory* svm_mem = amd::MemObjMap::FindMemObj(dev_ptr);
    if ((nullptr == svm_mem) || ((svm_mem->getMemFlags() & CL_MEM_ALLOC_HOST_PTR) == 0) ||
        // Validate the range of provided memory
        ((svm_mem->getSize() - (reinterpret_cast<const_address>(dev_ptr) -
          reinterpret_cast<address>(svm_mem->getSvmPtr()))) < count)) {
      LogPrintfError("SetSvmAttributes received unknown memory for update: %p!", dev_ptr);
      return false;
    }
  }
  if (info().hmmSupported_) {
    std::vector<hsa_amd_svm_attribute_pair_t> attr;

    switch (advice) {
      case amd::MemoryAdvice::SetReadMostly:
        attr.push_back({HSA_AMD_SVM_ATTRIB_READ_MOSTLY, true});
        break;
      case amd::MemoryAdvice::UnsetReadMostly:
        attr.push_back({HSA_AMD_SVM_ATTRIB_READ_MOSTLY, false});
        break;
      case amd::MemoryAdvice::SetPreferredLocation:
        if (use_cpu) {
          attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getCpuAgent().handle});
        } else {
          attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getBackendDevice().handle});
        }
        break;
      case amd::MemoryAdvice::UnsetPreferredLocation:
        // @note: 0 may cause a failure on old runtimes
        attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, 0});
        break;
      case amd::MemoryAdvice::SetAccessedBy: {
        const uint64_t attrib = (first_alloc) ? HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE :
                                                HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE_IN_PLACE;
        if (use_cpu) {
          attr.push_back({attrib, getCpuAgent().handle});
        } else {
          if (first_alloc) {
            // Provide access to all possible devices.
            //! @note: HMM should support automatic page table update with xnack enabled,
            //! but currently it doesn't and runtime explicitly enables access from all devices
            for (const auto dev : devices()) {
              // Skip null devices
              if (static_cast<Device*>(dev)->getBackendDevice().handle != 0) {
                attr.push_back({attrib, static_cast<Device*>(dev)->getBackendDevice().handle});
              }
            }
          } else {
            attr.push_back({attrib, getBackendDevice().handle});
          }
        }
        break;
      }
      case amd::MemoryAdvice::UnsetAccessedBy:
        // When unsetting we should use HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE for the agent
        attr.push_back({HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE, getBackendDevice().handle});
        break;
      case amd::MemoryAdvice::SetCoarseGrain:
        attr.push_back({HSA_AMD_SVM_ATTRIB_GLOBAL_FLAG, HSA_AMD_SVM_GLOBAL_FLAG_COARSE_GRAINED});
        break;
      case amd::MemoryAdvice::UnsetCoarseGrain:
        attr.push_back({HSA_AMD_SVM_ATTRIB_GLOBAL_FLAG, HSA_AMD_SVM_GLOBAL_FLAG_FINE_GRAINED});
        break;
      default:
        return false;
      break;
    }

    hsa_status_t status = hsa_amd_svm_attributes_set(const_cast<void*>(dev_ptr), count,
                                                    attr.data(), attr.size());
    if (status != HSA_STATUS_SUCCESS) {
      LogPrintfError("hsa_amd_svm_attributes_set() failed. Advice: %d, status: %d", advice, status);
      return false;
    }
  } else {
    LogWarning("hsa_amd_svm_attributes_set() is ignored, because no HMM support");
  }
  return true;
}

// ================================================================================================
bool Device::SetSvmAttributes(const void* dev_ptr, size_t count,
                              amd::MemoryAdvice advice, bool use_cpu) const {
  constexpr bool kFirstAlloc = false;
  return SetSvmAttributesInt(dev_ptr, count, advice, kFirstAlloc, use_cpu);
}

// ================================================================================================
bool Device::GetSvmAttributes(void** data, size_t* data_sizes, int* attributes,
                              size_t num_attributes, const void* dev_ptr, size_t count) const {
  if (settings().hmmFlags_ & Settings::Hmm::EnableSvmTracking) {
    amd::Memory* svm_mem = amd::MemObjMap::FindMemObj(dev_ptr);
    if ((nullptr == svm_mem) || ((svm_mem->getMemFlags() & CL_MEM_ALLOC_HOST_PTR) == 0) ||
        // Validate the range of provided memory
        ((svm_mem->getSize() - (reinterpret_cast<const_address>(dev_ptr) -
          reinterpret_cast<address>(svm_mem->getSvmPtr()))) < count)) {
      LogPrintfError("GetSvmAttributes received unknown memory %p for state!", dev_ptr);
      return false;
    }
  }

  hsa_amd_pointer_info_t ptr_info = {};
  for (size_t i = 0; i < num_attributes; ++i) {
    if (attributes[i] == amd::MemRangeAttribute::CoherencyMode) {
      ptr_info.size = sizeof(hsa_amd_pointer_info_t);
      // Query ptr type to see if it's a HMM allocation
      hsa_status_t status = hsa_amd_pointer_info(
        const_cast<void*>(dev_ptr), &ptr_info, nullptr, nullptr, nullptr);
      // The call should never fail in ROCR, but just check for an error and continue
      if (status != HSA_STATUS_SUCCESS) {
        LogError("hsa_amd_pointer_info() failed");
      }
      // Check if it's a legacy non-HMM allocation and update query
      if (ptr_info.type != HSA_EXT_POINTER_TYPE_UNKNOWN) {
        if (ptr_info.global_flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) {
          *reinterpret_cast<uint32_t*>(data[i]) = HSA_AMD_SVM_GLOBAL_FLAG_COARSE_GRAINED;
        } else if (ptr_info.global_flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED) {
          *reinterpret_cast<uint32_t*>(data[i]) = HSA_AMD_SVM_GLOBAL_FLAG_FINE_GRAINED;
        }
      }
    }
  }

  if (info().hmmSupported_) {
    uint32_t accessed_by = 0;
    std::vector<hsa_amd_svm_attribute_pair_t> attr;

    for (size_t i = 0; i < num_attributes; ++i) {
      switch (attributes[i]) {
        case amd::MemRangeAttribute::ReadMostly:
          attr.push_back({HSA_AMD_SVM_ATTRIB_READ_MOSTLY, 0});
          break;
        case amd::MemRangeAttribute::PreferredLocation:
          attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, 0});
          break;
        case amd::MemRangeAttribute::AccessedBy:
          accessed_by = attr.size();
          // Add all GPU devices into the query
          for (const auto agent : getGpuAgents()) {
            attr.push_back({HSA_AMD_SVM_ATTRIB_ACCESS_QUERY, agent.handle});
          }
          // Add CPU devices
          for (const auto agent_info : getCpuAgents()) {
            attr.push_back({HSA_AMD_SVM_ATTRIB_ACCESS_QUERY, agent_info.agent.handle});
          }
          accessed_by = attr.size() - accessed_by;
          break;
        case amd::MemRangeAttribute::LastPrefetchLocation:
          attr.push_back({HSA_AMD_SVM_ATTRIB_PREFETCH_LOCATION, 0});
          break;
        case amd::MemRangeAttribute::CoherencyMode:
          if (ptr_info.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
            attr.push_back({HSA_AMD_SVM_ATTRIB_GLOBAL_FLAG, 0});
          }
          break;
        default:
          return false;
        break;
      }
    }

    hsa_status_t status = hsa_amd_svm_attributes_get(const_cast<void*>(dev_ptr), count,
                                                    attr.data(), attr.size());
    if (status != HSA_STATUS_SUCCESS) {
      LogError("hsa_amd_svm_attributes_get() failed");
      return false;
    }

    uint32_t idx = 0;
    uint32_t rocr_attr = 0;
    for (size_t i = 0; i < num_attributes; ++i) {
      const auto& it = attr[rocr_attr];
      switch (attributes[i]) {
        case amd::MemRangeAttribute::ReadMostly:
          if (data_sizes[idx] != sizeof(uint32_t)) {
            return false;
          }
          // Cast ROCr value into the hip format
          *reinterpret_cast<uint32_t*>(data[idx]) =
              (static_cast<uint32_t>(it.value) > 0) ? true : false;
          ++rocr_attr;
          break;
        // The logic should be identical for the both queries
        case amd::MemRangeAttribute::PreferredLocation:
        case amd::MemRangeAttribute::LastPrefetchLocation:
          if (data_sizes[idx] != sizeof(uint32_t)) {
            return false;
          }
          *reinterpret_cast<int32_t*>(data[idx]) = static_cast<int32_t>(amd::InvalidDeviceId);
          // Find device agent returned by ROCr
          for (auto& device : devices()) {
            if (static_cast<Device*>(device)->getBackendDevice().handle == it.value) {
              *reinterpret_cast<uint32_t*>(data[idx]) = static_cast<uint32_t>(device->index());
            }
          }
          // Find CPU agent returned by ROCr
          for (auto& agent_info : getCpuAgents()) {
            if (agent_info.agent.handle == it.value) {
              *reinterpret_cast<int32_t*>(data[idx]) = static_cast<int32_t>(amd::CpuDeviceId);
            }
          }
          ++rocr_attr;
          break;
        case amd::MemRangeAttribute::AccessedBy: {
          uint32_t entry = 0;
          uint32_t device_count = data_sizes[idx] / 4;
          // Make sure it's multiple of 4
          if (data_sizes[idx] % 4 != 0) {
            return false;
          }
          for (uint32_t att = 0; att < accessed_by; ++att) {
            const auto& it = attr[rocr_attr + att];
            if (entry >= device_count) {
              // The size of the array is less than the amount of available devices
              break;
            }
            switch (it.attribute) {
              case HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE:
              case HSA_AMD_SVM_ATTRIB_AGENT_NO_ACCESS:
                break;
              case HSA_AMD_SVM_ATTRIB_AGENT_ACCESSIBLE_IN_PLACE:
                reinterpret_cast<int32_t*>(data[idx])[entry] =
                  static_cast<int32_t>(amd::InvalidDeviceId);
                // Find device agent returned by ROCr
                for (auto& device : devices()) {
                  if (static_cast<Device*>(device)->getBackendDevice().handle == it.value) {
                    reinterpret_cast<uint32_t*>(data[idx])[entry] =
                      static_cast<uint32_t>(device->index());
                  }
                }
                // Find CPU agent returned by ROCr
                for (auto& agent_info : getCpuAgents()) {
                  if (agent_info.agent.handle == it.value) {
                    reinterpret_cast<int32_t*>(data[idx])[entry] =
                      static_cast<int32_t>(amd::CpuDeviceId);
                  }
                }
                ++entry;
                break;
              default:
                LogWarning("Unexpected result from HSA_AMD_SVM_ATTRIB_ACCESS_QUERY");
                break;
            }
          }
          rocr_attr += accessed_by;
          for (uint32_t i = entry; i < device_count; ++i) {
            reinterpret_cast<int32_t*>(data[idx])[i] =
              static_cast<int32_t>(amd::InvalidDeviceId);
          }
          break;
        }
        case amd::MemRangeAttribute::CoherencyMode:
          if (data_sizes[idx] != sizeof(uint32_t)) {
            return false;
          }
          // if ptr is HMM alloc then overwrite the values
          if (ptr_info.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
            // Cast ROCr value into the hip format
            *reinterpret_cast<uint32_t*>(data[idx]) = static_cast<uint32_t>(it.value);
          }
          ++rocr_attr;
          break;
        default:
          return false;
        break;
      }
      // Find the next location in the query
      ++idx;
    }
  } else if (ptr_info.type == HSA_EXT_POINTER_TYPE_UNKNOWN) {
    LogError("GetSvmAttributes() failed, because no HMM support");
    return false;
  }

  return true;
}

// ================================================================================================
bool Device::SvmAllocInit(void* memory, size_t size) const {
  amd::MemoryAdvice advice = amd::MemoryAdvice::SetAccessedBy;
  constexpr bool kFirstAlloc = true;
  if (!SetSvmAttributesInt(memory, size, advice, kFirstAlloc)) {
    return false;
  }

  if ((settings().hmmFlags_ & Settings::Hmm::EnableMallocPrefetch) == 0) {
    return true;
  }

  if (info().hmmSupported_) {
    // Initialize signal for the barrier
    hsa_signal_store_relaxed(prefetch_signal_, kInitSignalValueOne);

    // Initiate a prefetch command which should force memory update in HMM
    hsa_status_t status = hsa_amd_svm_prefetch_async(memory, size, getBackendDevice(),
                                                     0, nullptr, prefetch_signal_);
    if (status != HSA_STATUS_SUCCESS) {
      LogError("hsa_amd_svm_prefetch_async() failed");
      return false;
    }

    // Wait for the prefetch
    if (!WaitForSignal(prefetch_signal_)) {
      LogError("Barrier packet submission failed");
      return false;
    }
  } else {
    LogWarning("Early prefetch failed, because no HMM support");
  }

  return true;
}

// ================================================================================================
void Device::svmFree(void* ptr) const {
  amd::Memory* svmMem = amd::MemObjMap::FindMemObj(ptr);
  if (nullptr != svmMem) {
    amd::MemObjMap::RemoveMemObj(svmMem->getSvmPtr());
    svmMem->release();
  }
}

// ================================================================================================
VirtualGPU* Device::xferQueue() const {
  if (!xferQueue_) {
    // Create virtual device for internal memory transfer
    Device* thisDevice = const_cast<Device*>(this);
    thisDevice->xferQueue_ = reinterpret_cast<VirtualGPU*>(thisDevice->createVirtualDevice());
    if (!xferQueue_) {
      LogError("Couldn't create the device transfer manager!");
      return nullptr;
    }
  }
  xferQueue_->enableSyncBlit();
  return xferQueue_;
}

// ================================================================================================
bool Device::SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput,
  cl_set_device_clock_mode_output_amd* pSetClockModeOutput) {
  bool result = true;
  return result;
}

// ================================================================================================
bool Device::IsHwEventReadyForcedWait(const amd::Event& event) const {
  void* hw_event =
      (event.NotifyEvent() != nullptr) ? event.NotifyEvent()->HwEvent() : event.HwEvent();
  if (hw_event == nullptr) {
    ClPrint(amd::LOG_INFO, amd::LOG_SIG, "No HW event");
    return false;
  }
  static constexpr bool Timeout = true;
  return WaitForSignal<Timeout>(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, false, true);
}

// ================================================================================================
bool Device::IsHwEventReady(const amd::Event& event, bool wait, uint32_t hip_event_flags) const {
  void* hw_event =
      (event.NotifyEvent() != nullptr) ? event.NotifyEvent()->HwEvent() : event.HwEvent();
  if (hw_event == nullptr) {
    ClPrint(amd::LOG_INFO, amd::LOG_SIG, "No HW event");
    return false;
  } else if (wait) {
    // hipEventBlockingSync
    // when set the CPU gives up host thread for other work
    // when not set the CPU enters a busy-wait on the event to occur
    constexpr int kHipEventBlockingSync = 0x1;
    bool active_wait = !(hip_event_flags & kHipEventBlockingSync) && ActiveWait();
    return WaitForSignal(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, active_wait);
  }
  return (hsa_signal_load_relaxed(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_) == 0);
}

// ================================================================================================
void Device::getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const {
  void* hw_event = (event.NotifyEvent() != nullptr) ?
    event.NotifyEvent()->HwEvent() : event.HwEvent();
  if (hw_event == nullptr) {
    ClPrint(amd::LOG_INFO, amd::LOG_SIG, "No HW event to read time");
    *start = *end = 0;
  } else {
    fetchSignalTime(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, getBackendDevice(),
                    start, end);
  }
}

// ================================================================================================
static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
  if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {
    // Abort on device exceptions.
    const char* errorMsg = 0;
    hsa_status_string(status, &errorMsg);
    if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
      size_t global_available_mem = 0;
      Device* dev = reinterpret_cast<Device*>(data);
      if (HSA_STATUS_SUCCESS != hsa_agent_get_info(dev->getBackendDevice(),
                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MEMORY_AVAIL),
                         &global_available_mem)) {
        LogError("HSA_AMD_AGENT_INFO_MEMORY_AVAIL query failed.");
      }
      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
              "Callback: Queue %p Aborting with error : %s Code: 0x%x Available Free mem : %zu MB",
              queue->base_address, errorMsg, status, global_available_mem/Mi);
    } else {
      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
        "Callback: Queue %p aborting with error : %s code: 0x%x", queue->base_address,
        errorMsg, status);
    }
    abort();
  }
}

// ================================================================================================
hsa_queue_t* Device::getQueueFromPool(const uint qIndex) {
  // Check if queue with refCount 0 is available to use
  if (queuePool_[qIndex].size() < GPU_MAX_HW_QUEUES) {
    for (auto it = queuePool_[qIndex].begin(); it != queuePool_[qIndex].end(); it++) {
      if (it->second.refCount == 0) {
        it->second.refCount++;
        ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Selected queue refCount: %p (%d)",
                it->first->base_address, it->second.refCount);
        return it->first;
      }
    }
  } else {
    if (qIndex < QueuePriority::Total && queuePool_[qIndex].size() > 0) {
      typedef decltype(queuePool_)::value_type::const_reference PoolRef;
      auto lowest = std::min_element(
          queuePool_[qIndex].begin(), queuePool_[qIndex].end(),
          [](PoolRef A, PoolRef B) { return A.second.refCount < B.second.refCount; });
      lowest->second.refCount++;
      ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Selected queue refCount: %p (%d)",
              lowest->first->base_address, lowest->second.refCount);
      return lowest->first;
    }
  }
  return nullptr;
}

hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue,
                                  const std::vector<uint32_t>& cuMask,
                                  amd::CommandQueue::Priority priority) {
  assert(queuePool_[QueuePriority::Low].size() <= GPU_MAX_HW_QUEUES ||
         queuePool_[QueuePriority::Normal].size() <= GPU_MAX_HW_QUEUES ||
         queuePool_[QueuePriority::High].size() <= GPU_MAX_HW_QUEUES);

  ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Number of allocated hardware queues with low priority: %d,"
      " with normal priority: %d, with high priority: %d, maximum per priority is: %d",
      queuePool_[QueuePriority::Low].size(),
      queuePool_[QueuePriority::Normal].size(),
      queuePool_[QueuePriority::High].size(), GPU_MAX_HW_QUEUES);

  hsa_amd_queue_priority_t queue_priority;
  uint qIndex;
  switch (priority) {
    case amd::CommandQueue::Priority::Low:
      queue_priority = HSA_AMD_QUEUE_PRIORITY_LOW;
      qIndex = QueuePriority::Low;
      break;
    case amd::CommandQueue::Priority::High:
      queue_priority = HSA_AMD_QUEUE_PRIORITY_HIGH;
      qIndex = QueuePriority::High;
      break;
    case amd::CommandQueue::Priority::Normal:
    case amd::CommandQueue::Priority::Medium:
    default:
      queue_priority = HSA_AMD_QUEUE_PRIORITY_NORMAL;
      qIndex = QueuePriority::Normal;
      break;
  }

  // If we have reached the max number of queues, reuse an existing queue with the matching queue priority,
  // choosing the one with the least number of users.
  // Note: Don't attempt to reuse the cooperative queue, since it's single per device
  if (!coop_queue && (cuMask.size() == 0) &&
      ((queuePool_[qIndex].size() == GPU_MAX_HW_QUEUES) || queuePool_[qIndex].size() > 0)) {
    hsa_queue_t* queue = getQueueFromPool(qIndex);
    if (queue != nullptr) {
      return queue;
    }
  }

  // Else create a new queue. This also includes the initial state where there
  // is no queue.
  uint32_t queue_max_packets = 0;
  if (HSA_STATUS_SUCCESS !=
      hsa_agent_get_info(bkendDevice_, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &queue_max_packets)) {
    DevLogError("Cannot get hsa agent info \n");
    return nullptr;
  }
  auto queue_size = (queue_max_packets < queue_size_hint) ? queue_max_packets : queue_size_hint;

  hsa_queue_t* queue;
  auto queue_type = HSA_QUEUE_TYPE_MULTI;

  // Enable cooperative queue for the device queue
  if (coop_queue) {
    queue_type = HSA_QUEUE_TYPE_COOPERATIVE;
  }

  while (hsa_queue_create(bkendDevice_, queue_size, queue_type, callbackQueue, this,
                          std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
                          &queue) != HSA_STATUS_SUCCESS) {
    queue_size >>= 1;
    if (queue_size < 64) {
      // if a queue with the same requested priority available from the pool, returns it here
      if (!coop_queue && (cuMask.size() == 0) && (queuePool_[qIndex].size() > 0)) {
        return getQueueFromPool(qIndex);
      }
      DevLogError("Device::acquireQueue: hsa_queue_create failed!");
      return nullptr;
    }
  }

  // default priority is normal so no need to set it again
  if (queue_priority != HSA_AMD_QUEUE_PRIORITY_NORMAL) {
    hsa_status_t st =  hsa_amd_queue_set_priority(queue, queue_priority);
    if (st != HSA_STATUS_SUCCESS) {
      DevLogError("Device::acquireQueue: hsa_amd_queue_set_priority failed!");
      hsa_queue_destroy(queue);
      return nullptr;
    }
  }

  ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Created SWq=%p to map on HWq=%p with "
          "size %d with priority %d, cooperative: %i",
          queue, queue->base_address, queue_size, queue_priority, coop_queue);

  hsa_amd_profiling_set_profiler_enabled(queue, 1);
  if (cuMask.size() != 0 || info_.globalCUMask_.size() != 0) {
    std::stringstream ss;
    ss << std::hex;
    std::vector<uint32_t> mask = {};

    // handle scenarios where cuMask (custom-defined), globalCUMask_ or both are valid and
    // fill the final mask which will be appiled to the current queue
    if (cuMask.size() != 0 && info_.globalCUMask_.size() == 0) {
      mask = cuMask;
    } else if (cuMask.size() != 0 && info_.globalCUMask_.size() != 0) {
      for (unsigned int i = 0; i < std::min(cuMask.size(), info_.globalCUMask_.size()); i++) {
        mask.push_back(cuMask[i] & info_.globalCUMask_[i]);
      }
      // check to make sure after ANDing cuMask (custom-defined) with global
      //CU mask, we have non-zero mask, oterwise just apply global CU mask
      bool zeroCUMask = true;
      for (auto m : mask) {
        if (m != 0) {
          zeroCUMask = false;
          break;
        }
      }
      if (zeroCUMask) {
        mask = info_.globalCUMask_;
      }
    } else {
      mask = info_.globalCUMask_;
    }


    for (int i = mask.size() - 1; i >= 0; i--) {
      ss << mask[i];
    }
    ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Setting CU mask 0x%s for hardware queue %p",
            ss.str().c_str(), queue->base_address);

    std::vector<uint32_t> final_mask = {};
    // hsa_amd_queue_cu_set_mask expects each bit in cuMask to represent each CU
    // For wgp mode: Each wgp consists of 2 CUs and CUs must be adjacent pairwise enabled
    // Convert each bit in the cuMask from wgp to cu by duplicating it
    if (settings().enableWgpMode_) {
      final_mask.resize(mask.size() * 2, 0);

      for (int i = 0; i < mask.size(); i++) {
        for (int j = 0; j < 16; j++) {
          // Convert least significant 16 bits
          if (((mask[i] >> j) & 0x1) == 0x1) {
            final_mask[2 * i] |= (0x3 << (2 * j));
          }

          // Convert most significant 16 bits
          if (((mask[i] >> (16 + j)) & 0x1) == 0x1) {
            final_mask[2 * i + 1] |= (0x3 << (2 * j));
          }
        }
      }
    } else {
      final_mask = mask;
    }

    hsa_status_t status = hsa_amd_queue_cu_set_mask(queue,
                          final_mask.size() * 32, final_mask.data());
    if (status != HSA_STATUS_SUCCESS) {
      DevLogError("Device::acquireQueue: hsa_amd_queue_cu_set_mask failed!");
      hsa_queue_destroy(queue);
      return nullptr;
    }
    if (cuMask.size() != 0) {
      // add queues with custom CU mask into their special pool to keep track
      // of mapping of these queues to their associated queueInfo (i.e., hostcall buffers)
      auto result = queueWithCUMaskPool_[qIndex].emplace(std::make_pair(queue, QueueInfo()));
      assert(result.second && "QueueInfo already exists");
      auto& qInfo = result.first->second;
      qInfo.refCount = 1;

      return queue;
    }
  }

  if (coop_queue) {
    // Skip queue recycling for cooperative queues, since it should be just one
    // per device.
    return queue;
  }
  auto result = queuePool_[qIndex].emplace(std::make_pair(queue, QueueInfo()));
  assert(result.second && "QueueInfo already exists");
  auto &qInfo = result.first->second;
  qInfo.refCount = 1;
  ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "acquireQueue refCount: %p (%d)",
          result.first->first->base_address, result.first->second.refCount);
  return queue;
}

void Device::releaseQueue(hsa_queue_t* queue, const std::vector<uint32_t>& cuMask) {
  for (auto& it : cuMask.size() == 0 ? queuePool_ : queueWithCUMaskPool_) {
    auto qIter = it.find(queue);
    if (qIter != it.end()) {
      auto &qInfo = qIter->second;
      assert(qInfo.refCount > 0);
      qInfo.refCount--;
      ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "releaseQueue refCount:%p (%d)",
              qIter->first->base_address, qIter->second.refCount);
      // hsa queues with cumask set are not being reused. Hence, if the app uses multiple
      // such queues it can cause memory leak and those must be destroyed here once the
      // refcount reaches 0.
      if ((!cuMask.empty()) && (qInfo.refCount == 0)) {
        if (qInfo.hostcallBuffer_) {
          ClPrint(amd::LOG_INFO, amd::LOG_QUEUE,
                  "Deleting hostcall buffer %p for hardware queue %p",
                  qInfo.hostcallBuffer_, qIter->first->base_address);
          amd::disableHostcalls(qInfo.hostcallBuffer_);
          context().svmFree(qInfo.hostcallBuffer_);
        }
        ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Deleting hardware queue %p with refCount 0",
                queue->base_address);
        qIter = it.erase(qIter);
        hsa_queue_destroy(queue);
      }
    }
  }
}

void* Device::getOrCreateHostcallBuffer(hsa_queue_t* queue, bool coop_queue,
                                        const std::vector<uint32_t>& cuMask) {
  decltype(queuePool_)::value_type::iterator qIter;

  if (!coop_queue) {
    for (auto &it : cuMask.size() == 0 ? queuePool_ : queueWithCUMaskPool_) {
      qIter = it.find(queue);
      if (qIter != it.end()) {
        break;
      }
    }
    if (cuMask.size() == 0) {
      assert(qIter != queuePool_[QueuePriority::High].end());
    } else {
      assert(qIter != queueWithCUMaskPool_[QueuePriority::High].end());
    }

    if (qIter->second.hostcallBuffer_) {
      return qIter->second.hostcallBuffer_;
    }
  } else {
    if (coopHostcallBuffer_) {
      return coopHostcallBuffer_;
    }
  }

  // The number of packets required in each buffer is at least equal to the
  // maximum number of waves supported by the device.
  auto wavesPerCu = info().maxThreadsPerCU_ / info().wavefrontWidth_;
  auto numPackets = info().maxComputeUnits_ * wavesPerCu;

  auto size = amd::getHostcallBufferSize(numPackets);
  auto align = amd::getHostcallBufferAlignment();

  void* buffer = context().svmAlloc(size, align, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS);
  if (!buffer) {
    ClPrint(amd::LOG_ERROR, amd::LOG_QUEUE,
            "Failed to create hostcall buffer for hardware queue %p", queue->base_address);
    return nullptr;
  }
  ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "Created hostcall buffer %p for hardware queue %p", buffer,
          queue->base_address);
  if (!coop_queue) {
    qIter->second.hostcallBuffer_ = buffer;
  } else {
    coopHostcallBuffer_ = buffer;
  }
  if (!amd::enableHostcalls(*this, buffer, numPackets)) {
    ClPrint(amd::LOG_ERROR, amd::LOG_QUEUE, "Failed to register hostcall buffer %p with listener",
            buffer);
    return nullptr;
  }
  return buffer;
}

bool Device::findLinkInfo(const amd::Device& other_device,
                          std::vector<LinkAttrType>* link_attrs) {
  return findLinkInfo((static_cast<const roc::Device*>(&other_device))->gpuvm_segment_,
                       link_attrs);
}

bool Device::findLinkInfo(const hsa_amd_memory_pool_t& pool,
                          std::vector<LinkAttrType>* link_attrs) {

  if ((!pool.handle) || (link_attrs == nullptr)) {
    return false;
  }

  // Retrieve the hops between 2 devices.
  int32_t hops = 0;
  hsa_status_t hsa_status = hsa_amd_agent_memory_pool_get_info(bkendDevice_, pool,
                            HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS, &hops);

  if (hsa_status != HSA_STATUS_SUCCESS) {
    DevLogPrintfError("Cannot get hops info, hsa failed with status: %d", hsa_status);
    return false;
  }

  if (hops < 0) {
    return false;
  }

  // The pool is on its agent
  if (hops == 0) {
    for (auto& link_attr : (*link_attrs)) {
      switch (link_attr.first) {
        case kLinkLinkType: {
          // No link, so type is meaningless,
          // caller should ignore it
          link_attr.second = -1;
          break;
        }
        case kLinkHopCount: {
          // no hop
          link_attr.second = 0;
          break;
        }
        case kLinkDistance: {
          // distance is zero, if no hops
          link_attr.second = 0;
          break;
        }
        case kLinkAtomicSupport: {
          // atomic support if its on the same agent
          link_attr.second = 1;
          break;
        }
        default: {
          DevLogPrintfError("Invalid LinkAttribute: %d ", link_attr.first);
          return false;
        }
      }
    }
    return true;
  }

  // Retrieve link info on the pool.
  std::vector<hsa_amd_memory_pool_link_info_t> link_info(hops);
  hsa_status = hsa_amd_agent_memory_pool_get_info(bkendDevice_, pool,
               HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO, link_info.data());

  if (hsa_status != HSA_STATUS_SUCCESS) {
    DevLogPrintfError("Cannot retrieve link info, hsa failed with status: %d", hsa_status);
    return false;
  }

  for (auto& link_attr : (*link_attrs)) {
    switch (link_attr.first) {
      case kLinkLinkType: {
        link_attr.second = static_cast<int32_t>(link_info[0].link_type);
        break;
      }
      case kLinkHopCount: {
        uint32_t distance = 0;
        // Because of Rocrs limitation hops is set to 1 always between two different devices
        // If Rocr Changes the behaviour revisit this logic
        for (size_t hop_idx = 0; hop_idx < static_cast<size_t>(hops); ++hop_idx) {
          distance += link_info[hop_idx].numa_distance;
        }
        uint32_t oneHopDistance
          = (link_info[0].link_type == HSA_AMD_LINK_INFO_TYPE_XGMI) ? 13 : 20;
        link_attr.second = static_cast<int32_t>(distance/oneHopDistance);
        break;
      }
      case kLinkDistance: {
        uint32_t distance = 0;
        // Sum of distances between hops
        for (size_t hop_idx = 0; hop_idx < static_cast<size_t>(hops); ++hop_idx) {
          distance += link_info[hop_idx].numa_distance;
        }
        link_attr.second = static_cast<int32_t>(distance);
        break;
      }
      case kLinkAtomicSupport: {
        // if either of the atomic is supported
        link_attr.second = static_cast<int32_t>(link_info[0].atomic_support_64bit
                                                || link_info[0].atomic_support_32bit);
        break;
      }
      default: {
        DevLogPrintfError("Invalid LinkAttribute: %d ", link_attr.first);
        return false;
      }
    }
  }

  return true;
}

// ================================================================================================
void Device::getGlobalCUMask(std::string cuMaskStr) {
  if (cuMaskStr.length() != 0) {
    std::string pre = cuMaskStr.substr(0, 2);
    if (pre.compare("0x") == 0 || pre.compare("0X") == 0) {
      cuMaskStr = cuMaskStr.substr(2, cuMaskStr.length());
    }

    int end = cuMaskStr.length();

    // the number of current physical CUs compressed in 4-bits
    size_t compPhysicalCUs = static_cast<size_t>((settings().enableWgpMode_ ?
           info_.maxComputeUnits_ * 2 : info_.maxComputeUnits_)/ 4);

    // the number of final available compute units after applying the requested CU mask
    uint32_t availCUs = 0;

    // read numCharToRead characters (8 or less) from the cuMask string each time, convert
    // it into hex, and store it into the globalCUMask_. If the length of the cuMask string
    // is more than the compressed physical available CUs, ignore the rest
    for (unsigned i = 0; i < std::min(cuMaskStr.length(), compPhysicalCUs); i += 8) {
      int numCharToRead = (i + 8 <= compPhysicalCUs) ? 8 : compPhysicalCUs - 8;
      std::string temp = cuMaskStr.substr(std::max(0, end - numCharToRead),
          std::min(numCharToRead, end));
      end -= numCharToRead;
      unsigned long ul = 0;
      try {
        ul = std::stoul(temp, 0, 16);
      } catch (const std::invalid_argument&) {
        info_.globalCUMask_ = {};
        availCUs = 0;
        break;
      }
      info_.globalCUMask_.push_back(static_cast<uint32_t>(ul));
      // count number of set bits in ul to find the number of active CUs
      // in each iteration
      while (ul) {
        ul &= (ul - 1);
        availCUs++;
      }
    }
    //update the maxComputeUnits_ based on the requested CU mask
    if (availCUs != 0 && availCUs < compPhysicalCUs * 4) {
      info_.maxComputeUnits_ = settings().enableWgpMode_ ?
      availCUs / 2 : availCUs;
    } else {
      info_.globalCUMask_ = {};
    }
  } else {
    info_.globalCUMask_ = {};
  }
}

// ================================================================================================
device::Signal* Device::createSignal() const {
  return new roc::Signal();
}

// ================================================================================================
amd::Memory* Device::GetArenaMemObj(const void* ptr, size_t& offset, size_t size) {
  // Only create arena_mem_object if CPU memory is accessible from HMM
  // or if runtime received an interop from another ROCr's client
  // Disable arena for XNACK
  hsa_amd_pointer_info_t ptr_info = {};
  ptr_info.size = sizeof(hsa_amd_pointer_info_t);
  if (!IsValidAllocation(ptr, size, &ptr_info)) {
    return nullptr;
  }

  if (arena_mem_obj_ == nullptr) {
    arena_mem_obj_ = new (context()) amd::ArenaMemory(context());
    if ((arena_mem_obj_ != nullptr) && !arena_mem_obj_->create(nullptr)) {
      LogError("Arena Memory Creation failed!");
      arena_mem_obj_->release();
      arena_mem_obj_ = nullptr;
    }
    if (arena_mem_obj_ == nullptr) {
      return arena_mem_obj_;
    }
  }

  // Calculate the offset of the pointer.
  const void* dev_ptr = reinterpret_cast<void*>(
      arena_mem_obj_->getDeviceMemory(*arena_mem_obj_->getContext().devices()[0])
          ->virtualAddress());
  offset = reinterpret_cast<size_t>(ptr) - reinterpret_cast<size_t>(dev_ptr);

  return arena_mem_obj_;
}

// ================================================================================================
void Device::ReleaseGlobalSignal(void* signal) const {
  if (signal != nullptr) {
    reinterpret_cast<ProfilingSignal*>(signal)->release();
  }
}

// ================================================================================================
bool Device::IsValidAllocation(const void* dev_ptr, size_t size, hsa_amd_pointer_info_t* ptr_info) {
  // Query ptr type to see if it's a HMM allocation
  hsa_status_t status =
      hsa_amd_pointer_info(const_cast<void*>(dev_ptr), ptr_info, nullptr, nullptr, nullptr);
  // The call should never fail in ROCR, but just check for an error and continue
  if (status != HSA_STATUS_SUCCESS) {
    LogError("hsa_amd_pointer_info() failed");
  }

  // Return false for pinned memory. A true return may result in a race because
  // ROCclr may attempt to do a pin/copy/unpin underneath in a multithreaded environment
  if (ptr_info->type == HSA_EXT_POINTER_TYPE_LOCKED) {
    return false;
  }

  if (ptr_info->type != HSA_EXT_POINTER_TYPE_UNKNOWN) {
    if ((size != 0) &&
        ((reinterpret_cast<const_address>(dev_ptr) -
          reinterpret_cast<const_address>(ptr_info->agentBaseAddress)) > size)) {
      return false;
    }
    return true;
  }
  return false;
}

// ================================================================================================
void Device::HiddenHeapAlloc(const VirtualGPU& gpu) {
  auto HeapAllocOnly = [this, &gpu]() -> bool {
    // Allocate initial heap for device memory allocator
    static constexpr size_t HeapBufferSize = 128 * Ki;
    heap_buffer_ = createMemory(HeapBufferSize);
    if (initial_heap_size_ != 0) {
      initial_heap_size_ = amd::alignUp(initial_heap_size_, 2 * Mi);
      initial_heap_buffer_ = createMemory(initial_heap_size_);
    }
    if (heap_buffer_ == nullptr) {
      LogError("Heap buffer allocation failed!");
      return false;
    }
    return true;
  };
  std::call_once(heap_allocated_, HeapAllocOnly);
}

// ================================================================================================
void Device::HiddenHeapInit(const VirtualGPU& gpu) {
  auto HeapZeroOut = [this, &gpu]() -> bool {
    static constexpr size_t HeapBufferSize = 128 * Ki;
    bool result = static_cast<const KernelBlitManager&>(gpu.blitMgr())
                      .initHeap(heap_buffer_, initial_heap_buffer_, HeapBufferSize,
                                initial_heap_size_ / (2 * Mi));

    return result;
  };
  std::call_once(heap_initialized_, HeapZeroOut);
}

// ================================================================================================
uint32_t Device::fetchSDMAMask(const device::BlitManager* handle, bool readEngine) const {
  uint32_t engine = 0;
  {
    amd::ScopedLock lock(vgpusAccess());
    for (auto it = engineAssignMap_.rbegin(); it != engineAssignMap_.rend(); ++it) {
      // If blitManager handle is in the map return the engine ID else
      // add to the map
      if (it->second == handle) {
        engine = it->first;
        break;
      } else if (it->second == 0) {
        it->second = handle;
        engine = it->first;
        break;
      }
    }
  }

  return (readEngine ? maxSdmaReadMask_ : maxSdmaWriteMask_) & engine;
}

// ================================================================================================
void Device::getSdmaRWMasks(uint32_t* readMask, uint32_t* writeMask) const {
  *readMask = maxSdmaReadMask_;
  *writeMask = maxSdmaWriteMask_;
}

// ================================================================================================
void Device::resetSDMAMask(const device::BlitManager* handle) const {
  amd::ScopedLock lock(vgpusAccess());

  for (auto& it : engineAssignMap_) {
    if (it.second == handle) {
      it.second = 0;
      break;
    }
  }
}

// ================================================================================================
ProfilingSignal::~ProfilingSignal() {
  if (signal_.handle != 0) {
    if (hsa_signal_load_relaxed(signal_) > 0) {
      LogError("Runtime shouldn't destroy a signal that is still busy!");
      if (hsa_signal_wait_scacquire(signal_, HSA_SIGNAL_CONDITION_LT, kInitSignalValueOne,
                                    kUnlimitedWait, HSA_WAIT_STATE_BLOCKED) != 0) {
      }
    }
    hsa_signal_destroy(signal_);
  }
}

#if defined(__clang__)
#if __has_feature(address_sanitizer)
device::UriLocator* Device::createUriLocator() const {
  return new roc::UriLocator();
}
#endif
#endif
} // namespace amd::roc
#endif  // WITHOUT_HSA_BACKEND
